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::lexer::PtxToken;
48use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
49use crate::r#type::common::*;
50
51pub mod section_0 {
52    use super::*;
53    use crate::r#type::instruction::red_async::section_0::*;
54
55    // ============================================================================
56    // Generated enum parsers
57    // ============================================================================
58
59    impl PtxParser for CompletionMechanism {
60        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
61            // Try MbarrierCompleteTxBytes
62            {
63                let saved_pos = stream.position();
64                if stream
65                    .expect_string(".mbarrier::complete_tx::bytes")
66                    .is_ok()
67                {
68                    return Ok(CompletionMechanism::MbarrierCompleteTxBytes);
69                }
70                stream.set_position(saved_pos);
71            }
72            let span = stream
73                .peek()
74                .map(|(_, s)| s.clone())
75                .unwrap_or(Span { start: 0, end: 0 });
76            let expected = &[".mbarrier::complete_tx::bytes"];
77            let found = stream
78                .peek()
79                .map(|(t, _)| format!("{:?}", t))
80                .unwrap_or_else(|_| "<end of input>".to_string());
81            Err(crate::parser::unexpected_value(span, expected, found))
82        }
83    }
84
85    impl PtxParser for Op {
86        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
87            // Try Inc
88            {
89                let saved_pos = stream.position();
90                if stream.expect_string(".inc").is_ok() {
91                    return Ok(Op::Inc);
92                }
93                stream.set_position(saved_pos);
94            }
95            let saved_pos = stream.position();
96            // Try Dec
97            {
98                let saved_pos = stream.position();
99                if stream.expect_string(".dec").is_ok() {
100                    return Ok(Op::Dec);
101                }
102                stream.set_position(saved_pos);
103            }
104            stream.set_position(saved_pos);
105            let span = stream
106                .peek()
107                .map(|(_, s)| s.clone())
108                .unwrap_or(Span { start: 0, end: 0 });
109            let expected = &[".inc", ".dec"];
110            let found = stream
111                .peek()
112                .map(|(t, _)| format!("{:?}", t))
113                .unwrap_or_else(|_| "<end of input>".to_string());
114            Err(crate::parser::unexpected_value(span, expected, found))
115        }
116    }
117
118    impl PtxParser for Scope {
119        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
120            // Try Cluster
121            {
122                let saved_pos = stream.position();
123                if stream.expect_string(".cluster").is_ok() {
124                    return Ok(Scope::Cluster);
125                }
126                stream.set_position(saved_pos);
127            }
128            let span = stream
129                .peek()
130                .map(|(_, s)| s.clone())
131                .unwrap_or(Span { start: 0, end: 0 });
132            let expected = &[".cluster"];
133            let found = stream
134                .peek()
135                .map(|(t, _)| format!("{:?}", t))
136                .unwrap_or_else(|_| "<end of input>".to_string());
137            Err(crate::parser::unexpected_value(span, expected, found))
138        }
139    }
140
141    impl PtxParser for Sem {
142        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
143            // Try Relaxed
144            {
145                let saved_pos = stream.position();
146                if stream.expect_string(".relaxed").is_ok() {
147                    return Ok(Sem::Relaxed);
148                }
149                stream.set_position(saved_pos);
150            }
151            let span = stream
152                .peek()
153                .map(|(_, s)| s.clone())
154                .unwrap_or(Span { start: 0, end: 0 });
155            let expected = &[".relaxed"];
156            let found = stream
157                .peek()
158                .map(|(t, _)| format!("{:?}", t))
159                .unwrap_or_else(|_| "<end of input>".to_string());
160            Err(crate::parser::unexpected_value(span, expected, found))
161        }
162    }
163
164    impl PtxParser for Ss {
165        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
166            // Try SharedCluster
167            {
168                let saved_pos = stream.position();
169                if stream.expect_string(".shared::cluster").is_ok() {
170                    return Ok(Ss::SharedCluster);
171                }
172                stream.set_position(saved_pos);
173            }
174            let span = stream
175                .peek()
176                .map(|(_, s)| s.clone())
177                .unwrap_or(Span { start: 0, end: 0 });
178            let expected = &[".shared::cluster"];
179            let found = stream
180                .peek()
181                .map(|(t, _)| format!("{:?}", t))
182                .unwrap_or_else(|_| "<end of input>".to_string());
183            Err(crate::parser::unexpected_value(span, expected, found))
184        }
185    }
186
187    impl PtxParser for Type {
188        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
189            // Try U32
190            {
191                let saved_pos = stream.position();
192                if stream.expect_string(".u32").is_ok() {
193                    return Ok(Type::U32);
194                }
195                stream.set_position(saved_pos);
196            }
197            let span = stream
198                .peek()
199                .map(|(_, s)| s.clone())
200                .unwrap_or(Span { start: 0, end: 0 });
201            let expected = &[".u32"];
202            let found = stream
203                .peek()
204                .map(|(t, _)| format!("{:?}", t))
205                .unwrap_or_else(|_| "<end of input>".to_string());
206            Err(crate::parser::unexpected_value(span, expected, found))
207        }
208    }
209
210    impl PtxParser for RedAsyncSemScopeSsCompletionMechanismOpType {
211        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
212            stream.expect_string("red")?;
213            stream.expect_string(".async")?;
214            let async_ = ();
215            stream.expect_complete()?;
216            let sem = Sem::parse(stream)?;
217            stream.expect_complete()?;
218            let scope = Scope::parse(stream)?;
219            stream.expect_complete()?;
220            let saved_pos = stream.position();
221            let ss = match Ss::parse(stream) {
222                Ok(val) => Some(val),
223                Err(_) => {
224                    stream.set_position(saved_pos);
225                    None
226                }
227            };
228            stream.expect_complete()?;
229            let completion_mechanism = CompletionMechanism::parse(stream)?;
230            stream.expect_complete()?;
231            let op = Op::parse(stream)?;
232            stream.expect_complete()?;
233            let type_ = Type::parse(stream)?;
234            stream.expect_complete()?;
235            let a = AddressOperand::parse(stream)?;
236            stream.expect_complete()?;
237            stream.expect(&PtxToken::Comma)?;
238            let b = GeneralOperand::parse(stream)?;
239            stream.expect_complete()?;
240            stream.expect(&PtxToken::Comma)?;
241            let mbar = AddressOperand::parse(stream)?;
242            stream.expect_complete()?;
243            stream.expect_complete()?;
244            stream.expect(&PtxToken::Semicolon)?;
245            Ok(RedAsyncSemScopeSsCompletionMechanismOpType {
246                async_,
247                sem,
248                scope,
249                ss,
250                completion_mechanism,
251                op,
252                type_,
253                a,
254                b,
255                mbar,
256            })
257        }
258    }
259}
260
261pub mod section_1 {
262    use super::*;
263    use crate::r#type::instruction::red_async::section_1::*;
264
265    // ============================================================================
266    // Generated enum parsers
267    // ============================================================================
268
269    impl PtxParser for CompletionMechanism {
270        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
271            // Try MbarrierCompleteTxBytes
272            {
273                let saved_pos = stream.position();
274                if stream
275                    .expect_string(".mbarrier::complete_tx::bytes")
276                    .is_ok()
277                {
278                    return Ok(CompletionMechanism::MbarrierCompleteTxBytes);
279                }
280                stream.set_position(saved_pos);
281            }
282            let span = stream
283                .peek()
284                .map(|(_, s)| s.clone())
285                .unwrap_or(Span { start: 0, end: 0 });
286            let expected = &[".mbarrier::complete_tx::bytes"];
287            let found = stream
288                .peek()
289                .map(|(t, _)| format!("{:?}", t))
290                .unwrap_or_else(|_| "<end of input>".to_string());
291            Err(crate::parser::unexpected_value(span, expected, found))
292        }
293    }
294
295    impl PtxParser for Op {
296        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
297            // Try Min
298            {
299                let saved_pos = stream.position();
300                if stream.expect_string(".min").is_ok() {
301                    return Ok(Op::Min);
302                }
303                stream.set_position(saved_pos);
304            }
305            let saved_pos = stream.position();
306            // Try Max
307            {
308                let saved_pos = stream.position();
309                if stream.expect_string(".max").is_ok() {
310                    return Ok(Op::Max);
311                }
312                stream.set_position(saved_pos);
313            }
314            stream.set_position(saved_pos);
315            let span = stream
316                .peek()
317                .map(|(_, s)| s.clone())
318                .unwrap_or(Span { start: 0, end: 0 });
319            let expected = &[".min", ".max"];
320            let found = stream
321                .peek()
322                .map(|(t, _)| format!("{:?}", t))
323                .unwrap_or_else(|_| "<end of input>".to_string());
324            Err(crate::parser::unexpected_value(span, expected, found))
325        }
326    }
327
328    impl PtxParser for Scope {
329        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
330            // Try Cluster
331            {
332                let saved_pos = stream.position();
333                if stream.expect_string(".cluster").is_ok() {
334                    return Ok(Scope::Cluster);
335                }
336                stream.set_position(saved_pos);
337            }
338            let span = stream
339                .peek()
340                .map(|(_, s)| s.clone())
341                .unwrap_or(Span { start: 0, end: 0 });
342            let expected = &[".cluster"];
343            let found = stream
344                .peek()
345                .map(|(t, _)| format!("{:?}", t))
346                .unwrap_or_else(|_| "<end of input>".to_string());
347            Err(crate::parser::unexpected_value(span, expected, found))
348        }
349    }
350
351    impl PtxParser for Sem {
352        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
353            // Try Relaxed
354            {
355                let saved_pos = stream.position();
356                if stream.expect_string(".relaxed").is_ok() {
357                    return Ok(Sem::Relaxed);
358                }
359                stream.set_position(saved_pos);
360            }
361            let span = stream
362                .peek()
363                .map(|(_, s)| s.clone())
364                .unwrap_or(Span { start: 0, end: 0 });
365            let expected = &[".relaxed"];
366            let found = stream
367                .peek()
368                .map(|(t, _)| format!("{:?}", t))
369                .unwrap_or_else(|_| "<end of input>".to_string());
370            Err(crate::parser::unexpected_value(span, expected, found))
371        }
372    }
373
374    impl PtxParser for Ss {
375        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
376            // Try SharedCluster
377            {
378                let saved_pos = stream.position();
379                if stream.expect_string(".shared::cluster").is_ok() {
380                    return Ok(Ss::SharedCluster);
381                }
382                stream.set_position(saved_pos);
383            }
384            let span = stream
385                .peek()
386                .map(|(_, s)| s.clone())
387                .unwrap_or(Span { start: 0, end: 0 });
388            let expected = &[".shared::cluster"];
389            let found = stream
390                .peek()
391                .map(|(t, _)| format!("{:?}", t))
392                .unwrap_or_else(|_| "<end of input>".to_string());
393            Err(crate::parser::unexpected_value(span, expected, found))
394        }
395    }
396
397    impl PtxParser for Type {
398        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
399            // Try U32
400            {
401                let saved_pos = stream.position();
402                if stream.expect_string(".u32").is_ok() {
403                    return Ok(Type::U32);
404                }
405                stream.set_position(saved_pos);
406            }
407            let saved_pos = stream.position();
408            // Try S32
409            {
410                let saved_pos = stream.position();
411                if stream.expect_string(".s32").is_ok() {
412                    return Ok(Type::S32);
413                }
414                stream.set_position(saved_pos);
415            }
416            stream.set_position(saved_pos);
417            let span = stream
418                .peek()
419                .map(|(_, s)| s.clone())
420                .unwrap_or(Span { start: 0, end: 0 });
421            let expected = &[".u32", ".s32"];
422            let found = stream
423                .peek()
424                .map(|(t, _)| format!("{:?}", t))
425                .unwrap_or_else(|_| "<end of input>".to_string());
426            Err(crate::parser::unexpected_value(span, expected, found))
427        }
428    }
429
430    impl PtxParser for RedAsyncSemScopeSsCompletionMechanismOpType1 {
431        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
432            stream.expect_string("red")?;
433            stream.expect_string(".async")?;
434            let async_ = ();
435            stream.expect_complete()?;
436            let sem = Sem::parse(stream)?;
437            stream.expect_complete()?;
438            let scope = Scope::parse(stream)?;
439            stream.expect_complete()?;
440            let saved_pos = stream.position();
441            let ss = match Ss::parse(stream) {
442                Ok(val) => Some(val),
443                Err(_) => {
444                    stream.set_position(saved_pos);
445                    None
446                }
447            };
448            stream.expect_complete()?;
449            let completion_mechanism = CompletionMechanism::parse(stream)?;
450            stream.expect_complete()?;
451            let op = Op::parse(stream)?;
452            stream.expect_complete()?;
453            let type_ = Type::parse(stream)?;
454            stream.expect_complete()?;
455            let a = AddressOperand::parse(stream)?;
456            stream.expect_complete()?;
457            stream.expect(&PtxToken::Comma)?;
458            let b = GeneralOperand::parse(stream)?;
459            stream.expect_complete()?;
460            stream.expect(&PtxToken::Comma)?;
461            let mbar = AddressOperand::parse(stream)?;
462            stream.expect_complete()?;
463            stream.expect_complete()?;
464            stream.expect(&PtxToken::Semicolon)?;
465            Ok(RedAsyncSemScopeSsCompletionMechanismOpType1 {
466                async_,
467                sem,
468                scope,
469                ss,
470                completion_mechanism,
471                op,
472                type_,
473                a,
474                b,
475                mbar,
476            })
477        }
478    }
479}
480
481pub mod section_2 {
482    use super::*;
483    use crate::r#type::instruction::red_async::section_2::*;
484
485    // ============================================================================
486    // Generated enum parsers
487    // ============================================================================
488
489    impl PtxParser for CompletionMechanism {
490        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
491            // Try MbarrierCompleteTxBytes
492            {
493                let saved_pos = stream.position();
494                if stream
495                    .expect_string(".mbarrier::complete_tx::bytes")
496                    .is_ok()
497                {
498                    return Ok(CompletionMechanism::MbarrierCompleteTxBytes);
499                }
500                stream.set_position(saved_pos);
501            }
502            let span = stream
503                .peek()
504                .map(|(_, s)| s.clone())
505                .unwrap_or(Span { start: 0, end: 0 });
506            let expected = &[".mbarrier::complete_tx::bytes"];
507            let found = stream
508                .peek()
509                .map(|(t, _)| format!("{:?}", t))
510                .unwrap_or_else(|_| "<end of input>".to_string());
511            Err(crate::parser::unexpected_value(span, expected, found))
512        }
513    }
514
515    impl PtxParser for Op {
516        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
517            // Try And
518            {
519                let saved_pos = stream.position();
520                if stream.expect_string(".and").is_ok() {
521                    return Ok(Op::And);
522                }
523                stream.set_position(saved_pos);
524            }
525            let saved_pos = stream.position();
526            // Try Xor
527            {
528                let saved_pos = stream.position();
529                if stream.expect_string(".xor").is_ok() {
530                    return Ok(Op::Xor);
531                }
532                stream.set_position(saved_pos);
533            }
534            stream.set_position(saved_pos);
535            let saved_pos = stream.position();
536            // Try Or
537            {
538                let saved_pos = stream.position();
539                if stream.expect_string(".or").is_ok() {
540                    return Ok(Op::Or);
541                }
542                stream.set_position(saved_pos);
543            }
544            stream.set_position(saved_pos);
545            let span = stream
546                .peek()
547                .map(|(_, s)| s.clone())
548                .unwrap_or(Span { start: 0, end: 0 });
549            let expected = &[".and", ".xor", ".or"];
550            let found = stream
551                .peek()
552                .map(|(t, _)| format!("{:?}", t))
553                .unwrap_or_else(|_| "<end of input>".to_string());
554            Err(crate::parser::unexpected_value(span, expected, found))
555        }
556    }
557
558    impl PtxParser for Scope {
559        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
560            // Try Cluster
561            {
562                let saved_pos = stream.position();
563                if stream.expect_string(".cluster").is_ok() {
564                    return Ok(Scope::Cluster);
565                }
566                stream.set_position(saved_pos);
567            }
568            let span = stream
569                .peek()
570                .map(|(_, s)| s.clone())
571                .unwrap_or(Span { start: 0, end: 0 });
572            let expected = &[".cluster"];
573            let found = stream
574                .peek()
575                .map(|(t, _)| format!("{:?}", t))
576                .unwrap_or_else(|_| "<end of input>".to_string());
577            Err(crate::parser::unexpected_value(span, expected, found))
578        }
579    }
580
581    impl PtxParser for Sem {
582        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
583            // Try Relaxed
584            {
585                let saved_pos = stream.position();
586                if stream.expect_string(".relaxed").is_ok() {
587                    return Ok(Sem::Relaxed);
588                }
589                stream.set_position(saved_pos);
590            }
591            let span = stream
592                .peek()
593                .map(|(_, s)| s.clone())
594                .unwrap_or(Span { start: 0, end: 0 });
595            let expected = &[".relaxed"];
596            let found = stream
597                .peek()
598                .map(|(t, _)| format!("{:?}", t))
599                .unwrap_or_else(|_| "<end of input>".to_string());
600            Err(crate::parser::unexpected_value(span, expected, found))
601        }
602    }
603
604    impl PtxParser for Ss {
605        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
606            // Try SharedCluster
607            {
608                let saved_pos = stream.position();
609                if stream.expect_string(".shared::cluster").is_ok() {
610                    return Ok(Ss::SharedCluster);
611                }
612                stream.set_position(saved_pos);
613            }
614            let span = stream
615                .peek()
616                .map(|(_, s)| s.clone())
617                .unwrap_or(Span { start: 0, end: 0 });
618            let expected = &[".shared::cluster"];
619            let found = stream
620                .peek()
621                .map(|(t, _)| format!("{:?}", t))
622                .unwrap_or_else(|_| "<end of input>".to_string());
623            Err(crate::parser::unexpected_value(span, expected, found))
624        }
625    }
626
627    impl PtxParser for Type {
628        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
629            // Try B32
630            {
631                let saved_pos = stream.position();
632                if stream.expect_string(".b32").is_ok() {
633                    return Ok(Type::B32);
634                }
635                stream.set_position(saved_pos);
636            }
637            let span = stream
638                .peek()
639                .map(|(_, s)| s.clone())
640                .unwrap_or(Span { start: 0, end: 0 });
641            let expected = &[".b32"];
642            let found = stream
643                .peek()
644                .map(|(t, _)| format!("{:?}", t))
645                .unwrap_or_else(|_| "<end of input>".to_string());
646            Err(crate::parser::unexpected_value(span, expected, found))
647        }
648    }
649
650    impl PtxParser for RedAsyncSemScopeSsCompletionMechanismOpType2 {
651        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
652            stream.expect_string("red")?;
653            stream.expect_string(".async")?;
654            let async_ = ();
655            stream.expect_complete()?;
656            let sem = Sem::parse(stream)?;
657            stream.expect_complete()?;
658            let scope = Scope::parse(stream)?;
659            stream.expect_complete()?;
660            let saved_pos = stream.position();
661            let ss = match Ss::parse(stream) {
662                Ok(val) => Some(val),
663                Err(_) => {
664                    stream.set_position(saved_pos);
665                    None
666                }
667            };
668            stream.expect_complete()?;
669            let completion_mechanism = CompletionMechanism::parse(stream)?;
670            stream.expect_complete()?;
671            let op = Op::parse(stream)?;
672            stream.expect_complete()?;
673            let type_ = Type::parse(stream)?;
674            stream.expect_complete()?;
675            let a = AddressOperand::parse(stream)?;
676            stream.expect_complete()?;
677            stream.expect(&PtxToken::Comma)?;
678            let b = GeneralOperand::parse(stream)?;
679            stream.expect_complete()?;
680            stream.expect(&PtxToken::Comma)?;
681            let mbar = AddressOperand::parse(stream)?;
682            stream.expect_complete()?;
683            stream.expect_complete()?;
684            stream.expect(&PtxToken::Semicolon)?;
685            Ok(RedAsyncSemScopeSsCompletionMechanismOpType2 {
686                async_,
687                sem,
688                scope,
689                ss,
690                completion_mechanism,
691                op,
692                type_,
693                a,
694                b,
695                mbar,
696            })
697        }
698    }
699}
700
701pub mod section_3 {
702    use super::*;
703    use crate::r#type::instruction::red_async::section_3::*;
704
705    // ============================================================================
706    // Generated enum parsers
707    // ============================================================================
708
709    impl PtxParser for CompletionMechanism {
710        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
711            // Try MbarrierCompleteTxBytes
712            {
713                let saved_pos = stream.position();
714                if stream
715                    .expect_string(".mbarrier::complete_tx::bytes")
716                    .is_ok()
717                {
718                    return Ok(CompletionMechanism::MbarrierCompleteTxBytes);
719                }
720                stream.set_position(saved_pos);
721            }
722            let span = stream
723                .peek()
724                .map(|(_, s)| s.clone())
725                .unwrap_or(Span { start: 0, end: 0 });
726            let expected = &[".mbarrier::complete_tx::bytes"];
727            let found = stream
728                .peek()
729                .map(|(t, _)| format!("{:?}", t))
730                .unwrap_or_else(|_| "<end of input>".to_string());
731            Err(crate::parser::unexpected_value(span, expected, found))
732        }
733    }
734
735    impl PtxParser for Scope {
736        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
737            // Try Cluster
738            {
739                let saved_pos = stream.position();
740                if stream.expect_string(".cluster").is_ok() {
741                    return Ok(Scope::Cluster);
742                }
743                stream.set_position(saved_pos);
744            }
745            let span = stream
746                .peek()
747                .map(|(_, s)| s.clone())
748                .unwrap_or(Span { start: 0, end: 0 });
749            let expected = &[".cluster"];
750            let found = stream
751                .peek()
752                .map(|(t, _)| format!("{:?}", t))
753                .unwrap_or_else(|_| "<end of input>".to_string());
754            Err(crate::parser::unexpected_value(span, expected, found))
755        }
756    }
757
758    impl PtxParser for Sem {
759        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
760            // Try Relaxed
761            {
762                let saved_pos = stream.position();
763                if stream.expect_string(".relaxed").is_ok() {
764                    return Ok(Sem::Relaxed);
765                }
766                stream.set_position(saved_pos);
767            }
768            let span = stream
769                .peek()
770                .map(|(_, s)| s.clone())
771                .unwrap_or(Span { start: 0, end: 0 });
772            let expected = &[".relaxed"];
773            let found = stream
774                .peek()
775                .map(|(t, _)| format!("{:?}", t))
776                .unwrap_or_else(|_| "<end of input>".to_string());
777            Err(crate::parser::unexpected_value(span, expected, found))
778        }
779    }
780
781    impl PtxParser for Ss {
782        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
783            // Try SharedCluster
784            {
785                let saved_pos = stream.position();
786                if stream.expect_string(".shared::cluster").is_ok() {
787                    return Ok(Ss::SharedCluster);
788                }
789                stream.set_position(saved_pos);
790            }
791            let span = stream
792                .peek()
793                .map(|(_, s)| s.clone())
794                .unwrap_or(Span { start: 0, end: 0 });
795            let expected = &[".shared::cluster"];
796            let found = stream
797                .peek()
798                .map(|(t, _)| format!("{:?}", t))
799                .unwrap_or_else(|_| "<end of input>".to_string());
800            Err(crate::parser::unexpected_value(span, expected, found))
801        }
802    }
803
804    impl PtxParser for Type {
805        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
806            // Try U32
807            {
808                let saved_pos = stream.position();
809                if stream.expect_string(".u32").is_ok() {
810                    return Ok(Type::U32);
811                }
812                stream.set_position(saved_pos);
813            }
814            let saved_pos = stream.position();
815            // Try S32
816            {
817                let saved_pos = stream.position();
818                if stream.expect_string(".s32").is_ok() {
819                    return Ok(Type::S32);
820                }
821                stream.set_position(saved_pos);
822            }
823            stream.set_position(saved_pos);
824            let saved_pos = stream.position();
825            // Try U64
826            {
827                let saved_pos = stream.position();
828                if stream.expect_string(".u64").is_ok() {
829                    return Ok(Type::U64);
830                }
831                stream.set_position(saved_pos);
832            }
833            stream.set_position(saved_pos);
834            let span = stream
835                .peek()
836                .map(|(_, s)| s.clone())
837                .unwrap_or(Span { start: 0, end: 0 });
838            let expected = &[".u32", ".s32", ".u64"];
839            let found = stream
840                .peek()
841                .map(|(t, _)| format!("{:?}", t))
842                .unwrap_or_else(|_| "<end of input>".to_string());
843            Err(crate::parser::unexpected_value(span, expected, found))
844        }
845    }
846
847    impl PtxParser for RedAsyncSemScopeSsCompletionMechanismAddType {
848        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
849            stream.expect_string("red")?;
850            stream.expect_string(".async")?;
851            let async_ = ();
852            stream.expect_complete()?;
853            let sem = Sem::parse(stream)?;
854            stream.expect_complete()?;
855            let scope = Scope::parse(stream)?;
856            stream.expect_complete()?;
857            let saved_pos = stream.position();
858            let ss = match Ss::parse(stream) {
859                Ok(val) => Some(val),
860                Err(_) => {
861                    stream.set_position(saved_pos);
862                    None
863                }
864            };
865            stream.expect_complete()?;
866            let completion_mechanism = CompletionMechanism::parse(stream)?;
867            stream.expect_complete()?;
868            stream.expect_string(".add")?;
869            let add = ();
870            stream.expect_complete()?;
871            let type_ = Type::parse(stream)?;
872            stream.expect_complete()?;
873            let a = AddressOperand::parse(stream)?;
874            stream.expect_complete()?;
875            stream.expect(&PtxToken::Comma)?;
876            let b = GeneralOperand::parse(stream)?;
877            stream.expect_complete()?;
878            stream.expect(&PtxToken::Comma)?;
879            let mbar = AddressOperand::parse(stream)?;
880            stream.expect_complete()?;
881            stream.expect_complete()?;
882            stream.expect(&PtxToken::Semicolon)?;
883            Ok(RedAsyncSemScopeSsCompletionMechanismAddType {
884                async_,
885                sem,
886                scope,
887                ss,
888                completion_mechanism,
889                add,
890                type_,
891                a,
892                b,
893                mbar,
894            })
895        }
896    }
897}
898
899pub mod section_4 {
900    use super::*;
901    use crate::r#type::instruction::red_async::section_4::*;
902
903    // ============================================================================
904    // Generated enum parsers
905    // ============================================================================
906
907    impl PtxParser for Scope {
908        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
909            // Try Cluster
910            {
911                let saved_pos = stream.position();
912                if stream.expect_string(".cluster").is_ok() {
913                    return Ok(Scope::Cluster);
914                }
915                stream.set_position(saved_pos);
916            }
917            let saved_pos = stream.position();
918            // Try Gpu
919            {
920                let saved_pos = stream.position();
921                if stream.expect_string(".gpu").is_ok() {
922                    return Ok(Scope::Gpu);
923                }
924                stream.set_position(saved_pos);
925            }
926            stream.set_position(saved_pos);
927            let span = stream
928                .peek()
929                .map(|(_, s)| s.clone())
930                .unwrap_or(Span { start: 0, end: 0 });
931            let expected = &[".cluster", ".gpu"];
932            let found = stream
933                .peek()
934                .map(|(t, _)| format!("{:?}", t))
935                .unwrap_or_else(|_| "<end of input>".to_string());
936            Err(crate::parser::unexpected_value(span, expected, found))
937        }
938    }
939
940    impl PtxParser for Sem {
941        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
942            // Try Release
943            {
944                let saved_pos = stream.position();
945                if stream.expect_string(".release").is_ok() {
946                    return Ok(Sem::Release);
947                }
948                stream.set_position(saved_pos);
949            }
950            let span = stream
951                .peek()
952                .map(|(_, s)| s.clone())
953                .unwrap_or(Span { start: 0, end: 0 });
954            let expected = &[".release"];
955            let found = stream
956                .peek()
957                .map(|(t, _)| format!("{:?}", t))
958                .unwrap_or_else(|_| "<end of input>".to_string());
959            Err(crate::parser::unexpected_value(span, expected, found))
960        }
961    }
962
963    impl PtxParser for Ss {
964        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
965            // Try Global
966            {
967                let saved_pos = stream.position();
968                if stream.expect_string(".global").is_ok() {
969                    return Ok(Ss::Global);
970                }
971                stream.set_position(saved_pos);
972            }
973            let span = stream
974                .peek()
975                .map(|(_, s)| s.clone())
976                .unwrap_or(Span { start: 0, end: 0 });
977            let expected = &[".global"];
978            let found = stream
979                .peek()
980                .map(|(t, _)| format!("{:?}", t))
981                .unwrap_or_else(|_| "<end of input>".to_string());
982            Err(crate::parser::unexpected_value(span, expected, found))
983        }
984    }
985
986    impl PtxParser for Type {
987        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
988            // Try U32
989            {
990                let saved_pos = stream.position();
991                if stream.expect_string(".u32").is_ok() {
992                    return Ok(Type::U32);
993                }
994                stream.set_position(saved_pos);
995            }
996            let saved_pos = stream.position();
997            // Try S32
998            {
999                let saved_pos = stream.position();
1000                if stream.expect_string(".s32").is_ok() {
1001                    return Ok(Type::S32);
1002                }
1003                stream.set_position(saved_pos);
1004            }
1005            stream.set_position(saved_pos);
1006            let saved_pos = stream.position();
1007            // Try U64
1008            {
1009                let saved_pos = stream.position();
1010                if stream.expect_string(".u64").is_ok() {
1011                    return Ok(Type::U64);
1012                }
1013                stream.set_position(saved_pos);
1014            }
1015            stream.set_position(saved_pos);
1016            let saved_pos = stream.position();
1017            // Try S64
1018            {
1019                let saved_pos = stream.position();
1020                if stream.expect_string(".s64").is_ok() {
1021                    return Ok(Type::S64);
1022                }
1023                stream.set_position(saved_pos);
1024            }
1025            stream.set_position(saved_pos);
1026            let span = stream
1027                .peek()
1028                .map(|(_, s)| s.clone())
1029                .unwrap_or(Span { start: 0, end: 0 });
1030            let expected = &[".u32", ".s32", ".u64", ".s64"];
1031            let found = stream
1032                .peek()
1033                .map(|(t, _)| format!("{:?}", t))
1034                .unwrap_or_else(|_| "<end of input>".to_string());
1035            Err(crate::parser::unexpected_value(span, expected, found))
1036        }
1037    }
1038
1039    impl PtxParser for RedAsyncMmioSemScopeSsAddType {
1040        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1041            stream.expect_string("red")?;
1042            stream.expect_string(".async")?;
1043            let async_ = ();
1044            stream.expect_complete()?;
1045            let saved_pos = stream.position();
1046            let mmio = stream.expect_string(".mmio").is_ok();
1047            if !mmio {
1048                stream.set_position(saved_pos);
1049            }
1050            stream.expect_complete()?;
1051            let sem = Sem::parse(stream)?;
1052            stream.expect_complete()?;
1053            let scope = Scope::parse(stream)?;
1054            stream.expect_complete()?;
1055            let saved_pos = stream.position();
1056            let ss = match Ss::parse(stream) {
1057                Ok(val) => Some(val),
1058                Err(_) => {
1059                    stream.set_position(saved_pos);
1060                    None
1061                }
1062            };
1063            stream.expect_complete()?;
1064            stream.expect_string(".add")?;
1065            let add = ();
1066            stream.expect_complete()?;
1067            let type_ = Type::parse(stream)?;
1068            stream.expect_complete()?;
1069            let a = AddressOperand::parse(stream)?;
1070            stream.expect_complete()?;
1071            stream.expect(&PtxToken::Comma)?;
1072            let b = GeneralOperand::parse(stream)?;
1073            stream.expect_complete()?;
1074            stream.expect_complete()?;
1075            stream.expect(&PtxToken::Semicolon)?;
1076            Ok(RedAsyncMmioSemScopeSsAddType {
1077                async_,
1078                mmio,
1079                sem,
1080                scope,
1081                ss,
1082                add,
1083                type_,
1084                a,
1085                b,
1086            })
1087        }
1088    }
1089}
ptx_parser/parser/instruction/red_async.rs

ptx_parser/parser/instruction/
red_async.rs
