ptx_parser/parser/instruction/
red_async.rs

1//! Original PTX specification:
2//!
3//! // Increment and Decrement reductions
4//! red.async.sem.scope{.ss}.completion_mechanism.op.type [a], b, [mbar];
5//! .sem  =                 { .relaxed };
6//! .scope =                { .cluster };
7//! .ss   =                 { .shared::cluster };
8//! .op   =                 { .inc, .dec };
9//! .type =                 { .u32 };
10//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
11//! ------------------------------------------------------------------
12//! // MIN and MAX reductions
13//! red.async.sem.scope{.ss}.completion_mechanism.op.type [a], b, [mbar];
14//! .sem  = { .relaxed };
15//! .scope = { .cluster };
16//! .ss   = { .shared::cluster };
17//! .op   = { .min, .max };
18//! .type = { .u32, .s32 };
19//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
20//! ------------------------------------------------------------------
21//! // Bitwise AND, OR and XOR reductions
22//! red.async.sem.scope{.ss}.completion_mechanism.op.type [a], b, [mbar];
23//! .sem  = { .relaxed };
24//! .scope = { .cluster };
25//! .ss   = { .shared::cluster };
26//! .op   = { .and, .or, .xor };
27//! .type = { .b32 };
28//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
29//! ------------------------------------------------------------------
30//! // ADD reductions
31//! red.async.sem.scope{.ss}.completion_mechanism.add.type [a], b, [mbar];
32//! .sem  = { .relaxed };
33//! .scope = { .cluster };
34//! .ss   = { .shared::cluster };
35//! .type = { .u32, .s32, .u64 };
36//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
37//! ----------------------------------------------------
38//! // Alternate floating point type:
39//! red.async{.mmio}.sem.scope{.ss}.add.type [a], b;
40//! .sem  = { .release };
41//! .scope = { .gpu, .cluster };
42//! .ss   = { .global };
43//! .type = { .u32, .s32, .u64, .s64 };
44
45#![allow(unused)]
46
47use crate::parser::{
48    PtxParseError, PtxParser, PtxTokenStream, Span,
49    util::{
50        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
51        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
52    },
53};
54use crate::r#type::common::*;
55use crate::{alt, ok, seq_n};
56
57pub mod section_0 {
58    use super::*;
59    use crate::r#type::instruction::red_async::section_0::*;
60
61    // ============================================================================
62    // Generated enum parsers
63    // ============================================================================
64
65    impl PtxParser for CompletionMechanism {
66        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
67            alt!(map(
68                string_p(".mbarrier::complete_tx::bytes"),
69                |_, _span| CompletionMechanism::MbarrierCompleteTxBytes
70            ))
71        }
72    }
73
74    impl PtxParser for Op {
75        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
76            alt!(
77                map(string_p(".inc"), |_, _span| Op::Inc),
78                map(string_p(".dec"), |_, _span| Op::Dec)
79            )
80        }
81    }
82
83    impl PtxParser for Scope {
84        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
85            alt!(map(string_p(".cluster"), |_, _span| Scope::Cluster))
86        }
87    }
88
89    impl PtxParser for Sem {
90        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
91            alt!(map(string_p(".relaxed"), |_, _span| Sem::Relaxed))
92        }
93    }
94
95    impl PtxParser for Ss {
96        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
97            alt!(map(string_p(".shared::cluster"), |_, _span| {
98                Ss::SharedCluster
99            }))
100        }
101    }
102
103    impl PtxParser for Type {
104        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
105            alt!(map(string_p(".u32"), |_, _span| Type::U32))
106        }
107    }
108
109    impl PtxParser for RedAsyncSemScopeSsCompletionMechanismOpType {
110        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
111            try_map(
112                seq_n!(
113                    string_p("red"),
114                    string_p(".async"),
115                    Sem::parse(),
116                    Scope::parse(),
117                    optional(Ss::parse()),
118                    CompletionMechanism::parse(),
119                    Op::parse(),
120                    Type::parse(),
121                    AddressOperand::parse(),
122                    comma_p(),
123                    GeneralOperand::parse(),
124                    comma_p(),
125                    AddressOperand::parse(),
126                    semicolon_p()
127                ),
128                |(
129                    _,
130                    async_,
131                    sem,
132                    scope,
133                    ss,
134                    completion_mechanism,
135                    op,
136                    type_,
137                    a,
138                    _,
139                    b,
140                    _,
141                    mbar,
142                    _,
143                ),
144                 span| {
145                    ok!(RedAsyncSemScopeSsCompletionMechanismOpType {
146                        async_ = async_,
147                        sem = sem,
148                        scope = scope,
149                        ss = ss,
150                        completion_mechanism = completion_mechanism,
151                        op = op,
152                        type_ = type_,
153                        a = a,
154                        b = b,
155                        mbar = mbar,
156
157                    })
158                },
159            )
160        }
161    }
162}
163
164pub mod section_1 {
165    use super::*;
166    use crate::r#type::instruction::red_async::section_1::*;
167
168    // ============================================================================
169    // Generated enum parsers
170    // ============================================================================
171
172    impl PtxParser for CompletionMechanism {
173        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
174            alt!(map(
175                string_p(".mbarrier::complete_tx::bytes"),
176                |_, _span| CompletionMechanism::MbarrierCompleteTxBytes
177            ))
178        }
179    }
180
181    impl PtxParser for Op {
182        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
183            alt!(
184                map(string_p(".min"), |_, _span| Op::Min),
185                map(string_p(".max"), |_, _span| Op::Max)
186            )
187        }
188    }
189
190    impl PtxParser for Scope {
191        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
192            alt!(map(string_p(".cluster"), |_, _span| Scope::Cluster))
193        }
194    }
195
196    impl PtxParser for Sem {
197        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
198            alt!(map(string_p(".relaxed"), |_, _span| Sem::Relaxed))
199        }
200    }
201
202    impl PtxParser for Ss {
203        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
204            alt!(map(string_p(".shared::cluster"), |_, _span| {
205                Ss::SharedCluster
206            }))
207        }
208    }
209
210    impl PtxParser for Type {
211        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
212            alt!(
213                map(string_p(".u32"), |_, _span| Type::U32),
214                map(string_p(".s32"), |_, _span| Type::S32)
215            )
216        }
217    }
218
219    impl PtxParser for RedAsyncSemScopeSsCompletionMechanismOpType1 {
220        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
221            try_map(
222                seq_n!(
223                    string_p("red"),
224                    string_p(".async"),
225                    Sem::parse(),
226                    Scope::parse(),
227                    optional(Ss::parse()),
228                    CompletionMechanism::parse(),
229                    Op::parse(),
230                    Type::parse(),
231                    AddressOperand::parse(),
232                    comma_p(),
233                    GeneralOperand::parse(),
234                    comma_p(),
235                    AddressOperand::parse(),
236                    semicolon_p()
237                ),
238                |(
239                    _,
240                    async_,
241                    sem,
242                    scope,
243                    ss,
244                    completion_mechanism,
245                    op,
246                    type_,
247                    a,
248                    _,
249                    b,
250                    _,
251                    mbar,
252                    _,
253                ),
254                 span| {
255                    ok!(RedAsyncSemScopeSsCompletionMechanismOpType1 {
256                        async_ = async_,
257                        sem = sem,
258                        scope = scope,
259                        ss = ss,
260                        completion_mechanism = completion_mechanism,
261                        op = op,
262                        type_ = type_,
263                        a = a,
264                        b = b,
265                        mbar = mbar,
266
267                    })
268                },
269            )
270        }
271    }
272}
273
274pub mod section_2 {
275    use super::*;
276    use crate::r#type::instruction::red_async::section_2::*;
277
278    // ============================================================================
279    // Generated enum parsers
280    // ============================================================================
281
282    impl PtxParser for CompletionMechanism {
283        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
284            alt!(map(
285                string_p(".mbarrier::complete_tx::bytes"),
286                |_, _span| CompletionMechanism::MbarrierCompleteTxBytes
287            ))
288        }
289    }
290
291    impl PtxParser for Op {
292        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
293            alt!(
294                map(string_p(".and"), |_, _span| Op::And),
295                map(string_p(".xor"), |_, _span| Op::Xor),
296                map(string_p(".or"), |_, _span| Op::Or)
297            )
298        }
299    }
300
301    impl PtxParser for Scope {
302        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
303            alt!(map(string_p(".cluster"), |_, _span| Scope::Cluster))
304        }
305    }
306
307    impl PtxParser for Sem {
308        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
309            alt!(map(string_p(".relaxed"), |_, _span| Sem::Relaxed))
310        }
311    }
312
313    impl PtxParser for Ss {
314        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
315            alt!(map(string_p(".shared::cluster"), |_, _span| {
316                Ss::SharedCluster
317            }))
318        }
319    }
320
321    impl PtxParser for Type {
322        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
323            alt!(map(string_p(".b32"), |_, _span| Type::B32))
324        }
325    }
326
327    impl PtxParser for RedAsyncSemScopeSsCompletionMechanismOpType2 {
328        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
329            try_map(
330                seq_n!(
331                    string_p("red"),
332                    string_p(".async"),
333                    Sem::parse(),
334                    Scope::parse(),
335                    optional(Ss::parse()),
336                    CompletionMechanism::parse(),
337                    Op::parse(),
338                    Type::parse(),
339                    AddressOperand::parse(),
340                    comma_p(),
341                    GeneralOperand::parse(),
342                    comma_p(),
343                    AddressOperand::parse(),
344                    semicolon_p()
345                ),
346                |(
347                    _,
348                    async_,
349                    sem,
350                    scope,
351                    ss,
352                    completion_mechanism,
353                    op,
354                    type_,
355                    a,
356                    _,
357                    b,
358                    _,
359                    mbar,
360                    _,
361                ),
362                 span| {
363                    ok!(RedAsyncSemScopeSsCompletionMechanismOpType2 {
364                        async_ = async_,
365                        sem = sem,
366                        scope = scope,
367                        ss = ss,
368                        completion_mechanism = completion_mechanism,
369                        op = op,
370                        type_ = type_,
371                        a = a,
372                        b = b,
373                        mbar = mbar,
374
375                    })
376                },
377            )
378        }
379    }
380}
381
382pub mod section_3 {
383    use super::*;
384    use crate::r#type::instruction::red_async::section_3::*;
385
386    // ============================================================================
387    // Generated enum parsers
388    // ============================================================================
389
390    impl PtxParser for CompletionMechanism {
391        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
392            alt!(map(
393                string_p(".mbarrier::complete_tx::bytes"),
394                |_, _span| CompletionMechanism::MbarrierCompleteTxBytes
395            ))
396        }
397    }
398
399    impl PtxParser for Scope {
400        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
401            alt!(map(string_p(".cluster"), |_, _span| Scope::Cluster))
402        }
403    }
404
405    impl PtxParser for Sem {
406        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
407            alt!(map(string_p(".relaxed"), |_, _span| Sem::Relaxed))
408        }
409    }
410
411    impl PtxParser for Ss {
412        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
413            alt!(map(string_p(".shared::cluster"), |_, _span| {
414                Ss::SharedCluster
415            }))
416        }
417    }
418
419    impl PtxParser for Type {
420        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
421            alt!(
422                map(string_p(".u32"), |_, _span| Type::U32),
423                map(string_p(".s32"), |_, _span| Type::S32),
424                map(string_p(".u64"), |_, _span| Type::U64)
425            )
426        }
427    }
428
429    impl PtxParser for RedAsyncSemScopeSsCompletionMechanismAddType {
430        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
431            try_map(
432                seq_n!(
433                    string_p("red"),
434                    string_p(".async"),
435                    Sem::parse(),
436                    Scope::parse(),
437                    optional(Ss::parse()),
438                    CompletionMechanism::parse(),
439                    string_p(".add"),
440                    Type::parse(),
441                    AddressOperand::parse(),
442                    comma_p(),
443                    GeneralOperand::parse(),
444                    comma_p(),
445                    AddressOperand::parse(),
446                    semicolon_p()
447                ),
448                |(
449                    _,
450                    async_,
451                    sem,
452                    scope,
453                    ss,
454                    completion_mechanism,
455                    add,
456                    type_,
457                    a,
458                    _,
459                    b,
460                    _,
461                    mbar,
462                    _,
463                ),
464                 span| {
465                    ok!(RedAsyncSemScopeSsCompletionMechanismAddType {
466                        async_ = async_,
467                        sem = sem,
468                        scope = scope,
469                        ss = ss,
470                        completion_mechanism = completion_mechanism,
471                        add = add,
472                        type_ = type_,
473                        a = a,
474                        b = b,
475                        mbar = mbar,
476
477                    })
478                },
479            )
480        }
481    }
482}
483
484pub mod section_4 {
485    use super::*;
486    use crate::r#type::instruction::red_async::section_4::*;
487
488    // ============================================================================
489    // Generated enum parsers
490    // ============================================================================
491
492    impl PtxParser for Scope {
493        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
494            alt!(
495                map(string_p(".cluster"), |_, _span| Scope::Cluster),
496                map(string_p(".gpu"), |_, _span| Scope::Gpu)
497            )
498        }
499    }
500
501    impl PtxParser for Sem {
502        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
503            alt!(map(string_p(".release"), |_, _span| Sem::Release))
504        }
505    }
506
507    impl PtxParser for Ss {
508        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
509            alt!(map(string_p(".global"), |_, _span| Ss::Global))
510        }
511    }
512
513    impl PtxParser for Type {
514        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
515            alt!(
516                map(string_p(".u32"), |_, _span| Type::U32),
517                map(string_p(".s32"), |_, _span| Type::S32),
518                map(string_p(".u64"), |_, _span| Type::U64),
519                map(string_p(".s64"), |_, _span| Type::S64)
520            )
521        }
522    }
523
524    impl PtxParser for RedAsyncMmioSemScopeSsAddType {
525        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
526            try_map(
527                seq_n!(
528                    string_p("red"),
529                    string_p(".async"),
530                    map(optional(string_p(".mmio")), |value, _| value.is_some()),
531                    Sem::parse(),
532                    Scope::parse(),
533                    optional(Ss::parse()),
534                    string_p(".add"),
535                    Type::parse(),
536                    AddressOperand::parse(),
537                    comma_p(),
538                    GeneralOperand::parse(),
539                    semicolon_p()
540                ),
541                |(_, async_, mmio, sem, scope, ss, add, type_, a, _, b, _), span| {
542                    ok!(RedAsyncMmioSemScopeSsAddType {
543                        async_ = async_,
544                        mmio = mmio,
545                        sem = sem,
546                        scope = scope,
547                        ss = ss,
548                        add = add,
549                        type_ = type_,
550                        a = a,
551                        b = b,
552
553                    })
554                },
555            )
556        }
557    }
558}
ptx_parser/parser/instruction/red_async.rs

ptx_parser/parser/instruction/
red_async.rs
