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

ptx_parser/parser/instruction/
red_async.rs
