ptx_parser/parser/instruction/
cp_reduce_async_bulk.rs

1//! Original PTX specification:
2//!
3//! cp.reduce.async.bulk.dst.src.completion_mechanism.redOp.type [dstMem], [srcMem], size, [mbar];
4//! .dst =                  { .shared::cluster };
5//! .src =                  { .shared::cta };
6//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
7//! .redOp=                 { .and, .or, .xor, .add, .inc, .dec, .min, .max };
8//! .type =                 { .b32, .u32, .s32, .b64, .u64 };
9//! ----------------------------------------------------------------
10//! cp.reduce.async.bulk.dst.src.completion_mechanism{.level::cache_hint}.redOp.type [dstMem], [srcMem], size{, cache-policy};
11//! .dst =                  { .global      };
12//! .src =                  { .shared::cta };
13//! ----------------------------------------------------------------
14//! .completion_mechanism = { .bulk_group };
15//! .level::cache_hint    = { .L2::cache_hint };
16//! .redOp=                 { .and, .or, .xor, .add, .inc, .dec, .min, .max };
17//! .type =                 { .f16, .bf16, .b32, .u32, .s32, .b64, .u64, .s64, .f32, .f64 };
18//! ----------------------------------------------------------------
19//! cp.reduce.async.bulk.dst.src.completion_mechanism{.level::cache_hint}.add.noftz.type [dstMem], [srcMem], size{, cache-policy};
20//! .dst  =                 { .global };
21//! .src  =                 { .shared::cta };
22//! .completion_mechanism = { .bulk_group };
23//! .type =                 { .f16, .bf16 };
24
25#![allow(unused)]
26
27use crate::parser::{
28    PtxParseError, PtxParser, PtxTokenStream, Span,
29    util::{
30        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
31        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
32    },
33};
34use crate::r#type::common::*;
35use crate::{alt, ok, seq_n};
36
37pub mod section_0 {
38    use super::*;
39    use crate::r#type::instruction::cp_reduce_async_bulk::section_0::*;
40
41    // ============================================================================
42    // Generated enum parsers
43    // ============================================================================
44
45    impl PtxParser for CompletionMechanism {
46        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
47            alt!(map(
48                string_p(".mbarrier::complete_tx::bytes"),
49                |_, _span| CompletionMechanism::MbarrierCompleteTxBytes
50            ))
51        }
52    }
53
54    impl PtxParser for Dst {
55        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
56            alt!(map(string_p(".shared::cluster"), |_, _span| {
57                Dst::SharedCluster
58            }))
59        }
60    }
61
62    impl PtxParser for Redop {
63        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
64            alt!(
65                map(string_p(".and"), |_, _span| Redop::And),
66                map(string_p(".xor"), |_, _span| Redop::Xor),
67                map(string_p(".add"), |_, _span| Redop::Add),
68                map(string_p(".inc"), |_, _span| Redop::Inc),
69                map(string_p(".dec"), |_, _span| Redop::Dec),
70                map(string_p(".min"), |_, _span| Redop::Min),
71                map(string_p(".max"), |_, _span| Redop::Max),
72                map(string_p(".or"), |_, _span| Redop::Or)
73            )
74        }
75    }
76
77    impl PtxParser for Src {
78        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
79            alt!(map(string_p(".shared::cta"), |_, _span| Src::SharedCta))
80        }
81    }
82
83    impl PtxParser for Type {
84        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
85            alt!(
86                map(string_p(".b32"), |_, _span| Type::B32),
87                map(string_p(".u32"), |_, _span| Type::U32),
88                map(string_p(".s32"), |_, _span| Type::S32),
89                map(string_p(".b64"), |_, _span| Type::B64),
90                map(string_p(".u64"), |_, _span| Type::U64)
91            )
92        }
93    }
94
95    impl PtxParser for CpReduceAsyncBulkDstSrcCompletionMechanismRedopType {
96        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
97            try_map(
98                seq_n!(
99                    string_p("cp"),
100                    string_p(".reduce"),
101                    string_p(".async"),
102                    string_p(".bulk"),
103                    Dst::parse(),
104                    Src::parse(),
105                    CompletionMechanism::parse(),
106                    Redop::parse(),
107                    Type::parse(),
108                    AddressOperand::parse(),
109                    comma_p(),
110                    AddressOperand::parse(),
111                    comma_p(),
112                    GeneralOperand::parse(),
113                    comma_p(),
114                    AddressOperand::parse(),
115                    semicolon_p()
116                ),
117                |(
118                    _,
119                    reduce,
120                    async_,
121                    bulk,
122                    dst,
123                    src,
124                    completion_mechanism,
125                    redop,
126                    type_,
127                    dstmem,
128                    _,
129                    srcmem,
130                    _,
131                    size,
132                    _,
133                    mbar,
134                    _,
135                ),
136                 span| {
137                    ok!(CpReduceAsyncBulkDstSrcCompletionMechanismRedopType {
138                        reduce = reduce,
139                        async_ = async_,
140                        bulk = bulk,
141                        dst = dst,
142                        src = src,
143                        completion_mechanism = completion_mechanism,
144                        redop = redop,
145                        type_ = type_,
146                        dstmem = dstmem,
147                        srcmem = srcmem,
148                        size = size,
149                        mbar = mbar,
150
151                    })
152                },
153            )
154        }
155    }
156}
157
158pub mod section_1 {
159    use super::*;
160    use crate::r#type::instruction::cp_reduce_async_bulk::section_1::*;
161
162    // ============================================================================
163    // Generated enum parsers
164    // ============================================================================
165
166    impl PtxParser for CompletionMechanism {
167        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
168            alt!(map(
169                string_p(".mbarrier::complete_tx::bytes"),
170                |_, _span| CompletionMechanism::MbarrierCompleteTxBytes
171            ))
172        }
173    }
174
175    impl PtxParser for Dst {
176        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
177            alt!(map(string_p(".global"), |_, _span| Dst::Global))
178        }
179    }
180
181    impl PtxParser for Redop {
182        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
183            alt!(
184                map(string_p(".and"), |_, _span| Redop::And),
185                map(string_p(".xor"), |_, _span| Redop::Xor),
186                map(string_p(".add"), |_, _span| Redop::Add),
187                map(string_p(".inc"), |_, _span| Redop::Inc),
188                map(string_p(".dec"), |_, _span| Redop::Dec),
189                map(string_p(".min"), |_, _span| Redop::Min),
190                map(string_p(".max"), |_, _span| Redop::Max),
191                map(string_p(".or"), |_, _span| Redop::Or)
192            )
193        }
194    }
195
196    impl PtxParser for Src {
197        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
198            alt!(map(string_p(".shared::cta"), |_, _span| Src::SharedCta))
199        }
200    }
201
202    impl PtxParser for Type {
203        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
204            alt!(
205                map(string_p(".b32"), |_, _span| Type::B32),
206                map(string_p(".u32"), |_, _span| Type::U32),
207                map(string_p(".s32"), |_, _span| Type::S32),
208                map(string_p(".b64"), |_, _span| Type::B64),
209                map(string_p(".u64"), |_, _span| Type::U64)
210            )
211        }
212    }
213
214    impl PtxParser for CpReduceAsyncBulkDstSrcCompletionMechanismLevelCacheHintRedopType {
215        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
216            try_map(
217                seq_n!(
218                    string_p("cp"),
219                    string_p(".reduce"),
220                    string_p(".async"),
221                    string_p(".bulk"),
222                    Dst::parse(),
223                    Src::parse(),
224                    CompletionMechanism::parse(),
225                    map(optional(string_p(".level::cache_hint")), |value, _| value
226                        .is_some()),
227                    Redop::parse(),
228                    Type::parse(),
229                    AddressOperand::parse(),
230                    comma_p(),
231                    AddressOperand::parse(),
232                    comma_p(),
233                    GeneralOperand::parse(),
234                    map(
235                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
236                        |value, _| value.map(|(_, operand)| operand)
237                    ),
238                    semicolon_p()
239                ),
240                |(
241                    _,
242                    reduce,
243                    async_,
244                    bulk,
245                    dst,
246                    src,
247                    completion_mechanism,
248                    level_cache_hint,
249                    redop,
250                    type_,
251                    dstmem,
252                    _,
253                    srcmem,
254                    _,
255                    size,
256                    cache_policy,
257                    _,
258                ),
259                 span| {
260                    ok!(CpReduceAsyncBulkDstSrcCompletionMechanismLevelCacheHintRedopType {
261                        reduce = reduce,
262                        async_ = async_,
263                        bulk = bulk,
264                        dst = dst,
265                        src = src,
266                        completion_mechanism = completion_mechanism,
267                        level_cache_hint = level_cache_hint,
268                        redop = redop,
269                        type_ = type_,
270                        dstmem = dstmem,
271                        srcmem = srcmem,
272                        size = size,
273                        cache_policy = cache_policy,
274
275                    })
276                },
277            )
278        }
279    }
280}
281
282pub mod section_2 {
283    use super::*;
284    use crate::r#type::instruction::cp_reduce_async_bulk::section_2::*;
285
286    // ============================================================================
287    // Generated enum parsers
288    // ============================================================================
289
290    impl PtxParser for CompletionMechanism {
291        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
292            alt!(map(string_p(".bulk_group"), |_, _span| {
293                CompletionMechanism::BulkGroup
294            }))
295        }
296    }
297
298    impl PtxParser for Dst {
299        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
300            alt!(map(string_p(".global"), |_, _span| Dst::Global))
301        }
302    }
303
304    impl PtxParser for LevelCacheHint {
305        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
306            alt!(map(string_p(".L2::cache_hint"), |_, _span| {
307                LevelCacheHint::L2CacheHint
308            }))
309        }
310    }
311
312    impl PtxParser for Src {
313        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
314            alt!(map(string_p(".shared::cta"), |_, _span| Src::SharedCta))
315        }
316    }
317
318    impl PtxParser for Type {
319        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
320            alt!(
321                map(string_p(".bf16"), |_, _span| Type::Bf16),
322                map(string_p(".f16"), |_, _span| Type::F16)
323            )
324        }
325    }
326
327    impl PtxParser for CpReduceAsyncBulkDstSrcCompletionMechanismLevelCacheHintAddNoftzType {
328        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
329            try_map(
330                seq_n!(
331                    string_p("cp"),
332                    string_p(".reduce"),
333                    string_p(".async"),
334                    string_p(".bulk"),
335                    Dst::parse(),
336                    Src::parse(),
337                    CompletionMechanism::parse(),
338                    optional(LevelCacheHint::parse()),
339                    string_p(".add"),
340                    string_p(".noftz"),
341                    Type::parse(),
342                    AddressOperand::parse(),
343                    comma_p(),
344                    AddressOperand::parse(),
345                    comma_p(),
346                    GeneralOperand::parse(),
347                    map(
348                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
349                        |value, _| value.map(|(_, operand)| operand)
350                    ),
351                    semicolon_p()
352                ),
353                |(
354                    _,
355                    reduce,
356                    async_,
357                    bulk,
358                    dst,
359                    src,
360                    completion_mechanism,
361                    level_cache_hint,
362                    add,
363                    noftz,
364                    type_,
365                    dstmem,
366                    _,
367                    srcmem,
368                    _,
369                    size,
370                    cache_policy,
371                    _,
372                ),
373                 span| {
374                    ok!(CpReduceAsyncBulkDstSrcCompletionMechanismLevelCacheHintAddNoftzType {
375                        reduce = reduce,
376                        async_ = async_,
377                        bulk = bulk,
378                        dst = dst,
379                        src = src,
380                        completion_mechanism = completion_mechanism,
381                        level_cache_hint = level_cache_hint,
382                        add = add,
383                        noftz = noftz,
384                        type_ = type_,
385                        dstmem = dstmem,
386                        srcmem = srcmem,
387                        size = size,
388                        cache_policy = cache_policy,
389
390                    })
391                },
392            )
393        }
394    }
395}
ptx_parser/parser/instruction/cp_reduce_async_bulk.rs

ptx_parser/parser/instruction/
cp_reduce_async_bulk.rs
