ptx_parser/parser/instruction/
mma.rs

1//! Original PTX specification:
2//!
3//! // Half precision floating point type:
4//! mma.sync.aligned.m8n8k4.alayout.blayout.dtype.f16.f16.ctype  d, a, b, c;
5//! mma.sync.aligned.m16n8k8.row.col.dtype.f16.f16.ctype  d, a, b, c;
6//! mma.sync.aligned.m16n8k16.row.col.dtype.f16.f16.ctype d, a, b, c;
7//! .alayout = {.row, .col};
8//! .blayout = {.row, .col};
9//! .ctype   = {.f16, .f32};
10//! .dtype   = {.f16, .f32};
11//! ----------------------------------------------------
12//! // Alternate floating point type:
13//! // Alternate floating point type:
14//! mma.sync.aligned.m16n8k4.row.col.f32.tf32.tf32.f32        d, a, b, c;
15//! mma.sync.aligned.m16n8k8.row.col.f32.atype.btype.f32      d, a, b, c;
16//! mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32       d, a, b, c;
17//! mma.sync.aligned.shape.row.col.dtype.f8type.f8type.ctype  d, a, b, c;
18//! mma.sync.aligned.m16n8k32.row.col.kind.dtype.f8f6f4type.f8f6f4type.ctype d, a, b, c;
19//! .atype      = {.bf16, .tf32};
20//! .btype      = {.bf16, .tf32};
21//! .f8type     = {.e4m3, .e5m2};
22//! .f8f6f4type = {.e4m3, .e5m2, .e3m2, .e2m3, .e2m1};
23//! .ctype      = {.f16, .f32};
24//! .dtype      = {.f16, .f32};
25//! .shape      = {.m16n8k16, .m16n8k32};
26//! .kind       = {.kind::f8f6f4};
27//! ----------------------------------------------------
28//! // Alternate floating point type:
29//! // Alternate floating point type with block scaling:
30//! mma.sync.aligned.m16n8k64.row.col.kind.block_scale{.scale_vec_size}.f32.e2m1.e2m1.f32.stype d, a, b, c, scale-a-data, {byte-id-a, thread-id-a}, scale-b-data, {byte-id-b, thread-id-b};
31//! .kind           = {.kind::mxf4};
32//! .scale_vec_size = {.scale_vec::2X};
33//! .stype          = {.ue8m0};
34//! ----------------------------------------------------
35//! // Alternate floating point type:
36//! mma.sync.aligned.m16n8k64.row.col.kind.block_scale.scale_vec_size.f32.e2m1.e2m1.f32.stype d, a, b, c, scale-a-data, {byte-id-a, thread-id-a}, scale-b-data, {byte-id-b, thread-id-b};
37//! .kind           = {.kind::mxf4nvf4};
38//! .scale_vec_size = {.scale_vec::2X, .scale_vec::4X};
39//! .stype          = {.ue8m0, .ue4m3};
40//! ----------------------------------------------------
41//! // Alternate floating point type:
42//! mma.sync.aligned.m16n8k32.row.col.kind.block_scale{.scale_vec_size}.f32.f8f6f4type.f8f6f4type.f32.stype d, a, b, c, scale-a-data, {byte-id-a, thread-id-a}, scale-b-data, {byte-id-b, thread-id-b};
43//! .kind           = {.kind::mxf8f6f4};
44//! .scale_vec_size = {.scale_vec::1X};
45//! .f8f6f4type     = {.e4m3, .e5m2, .e3m2, .e2m3, .e2m1};
46//! .stype          = {.ue8m0};
47//! ----------------------------------------------------
48//! // Alternate floating point type:
49//! // Double precision floating point type:
50//! mma.sync.aligned.shape.row.col.f64.f64.f64.f64 d, a, b, c;
51//! .shape   = {.m8n84, .m16n8k4, .m16n8k8, .m16n8k16};
52//! ----------------------------------------------------
53//! // Alternate floating point type:
54//! // Integer type:
55//! mma.sync.aligned.shape.row.col{.satfinite}.s32.atype.btype.s32 d, a, b, c;
56//! .shape   = {.m8n8k16, .m16n8k16, .m16n8k32};
57//! .atype   = {.u8, .s8};
58//! .btype   = {.u8, .s8};
59//! ----------------------------------------------------
60//! // Alternate floating point type:
61//! mma.sync.aligned.shape.row.col{.satfinite}.s32.atype.btype.s32 d, a, b, c;
62//! .shape   = {.m8n8k32, .m16n8k32, .m16n8k64};
63//! .atype   = {.u4, .s4};
64//! .btype   = {.u4, .s4};
65//! ----------------------------------------------------
66//! // Alternate floating point type:
67//! // Single bit:
68//! mma.sync.aligned.shape.row.col.s32.b1.b1.s32.bitOp.popc d, a, b, c;
69//! .bitOp = {.xor, .and};
70//! .shape = {.m8n8k128, .m16n8k128, .m16n8k256};
71
72#![allow(unused)]
73
74use crate::lexer::PtxToken;
75use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
76use crate::r#type::common::*;
77
78pub mod section_0 {
79    use super::*;
80    use crate::r#type::instruction::mma::section_0::*;
81
82    // ============================================================================
83    // Generated enum parsers
84    // ============================================================================
85
86    impl PtxParser for Alayout {
87        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
88            // Try Row
89            {
90                let saved_pos = stream.position();
91                if stream.expect_string(".row").is_ok() {
92                    return Ok(Alayout::Row);
93                }
94                stream.set_position(saved_pos);
95            }
96            let saved_pos = stream.position();
97            // Try Col
98            {
99                let saved_pos = stream.position();
100                if stream.expect_string(".col").is_ok() {
101                    return Ok(Alayout::Col);
102                }
103                stream.set_position(saved_pos);
104            }
105            stream.set_position(saved_pos);
106            let span = stream
107                .peek()
108                .map(|(_, s)| s.clone())
109                .unwrap_or(Span { start: 0, end: 0 });
110            let expected = &[".row", ".col"];
111            let found = stream
112                .peek()
113                .map(|(t, _)| format!("{:?}", t))
114                .unwrap_or_else(|_| "<end of input>".to_string());
115            Err(crate::parser::unexpected_value(span, expected, found))
116        }
117    }
118
119    impl PtxParser for Blayout {
120        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
121            // Try Row
122            {
123                let saved_pos = stream.position();
124                if stream.expect_string(".row").is_ok() {
125                    return Ok(Blayout::Row);
126                }
127                stream.set_position(saved_pos);
128            }
129            let saved_pos = stream.position();
130            // Try Col
131            {
132                let saved_pos = stream.position();
133                if stream.expect_string(".col").is_ok() {
134                    return Ok(Blayout::Col);
135                }
136                stream.set_position(saved_pos);
137            }
138            stream.set_position(saved_pos);
139            let span = stream
140                .peek()
141                .map(|(_, s)| s.clone())
142                .unwrap_or(Span { start: 0, end: 0 });
143            let expected = &[".row", ".col"];
144            let found = stream
145                .peek()
146                .map(|(t, _)| format!("{:?}", t))
147                .unwrap_or_else(|_| "<end of input>".to_string());
148            Err(crate::parser::unexpected_value(span, expected, found))
149        }
150    }
151
152    impl PtxParser for Ctype {
153        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
154            // Try F16
155            {
156                let saved_pos = stream.position();
157                if stream.expect_string(".f16").is_ok() {
158                    return Ok(Ctype::F16);
159                }
160                stream.set_position(saved_pos);
161            }
162            let saved_pos = stream.position();
163            // Try F32
164            {
165                let saved_pos = stream.position();
166                if stream.expect_string(".f32").is_ok() {
167                    return Ok(Ctype::F32);
168                }
169                stream.set_position(saved_pos);
170            }
171            stream.set_position(saved_pos);
172            let span = stream
173                .peek()
174                .map(|(_, s)| s.clone())
175                .unwrap_or(Span { start: 0, end: 0 });
176            let expected = &[".f16", ".f32"];
177            let found = stream
178                .peek()
179                .map(|(t, _)| format!("{:?}", t))
180                .unwrap_or_else(|_| "<end of input>".to_string());
181            Err(crate::parser::unexpected_value(span, expected, found))
182        }
183    }
184
185    impl PtxParser for Dtype {
186        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
187            // Try F16
188            {
189                let saved_pos = stream.position();
190                if stream.expect_string(".f16").is_ok() {
191                    return Ok(Dtype::F16);
192                }
193                stream.set_position(saved_pos);
194            }
195            let saved_pos = stream.position();
196            // Try F32
197            {
198                let saved_pos = stream.position();
199                if stream.expect_string(".f32").is_ok() {
200                    return Ok(Dtype::F32);
201                }
202                stream.set_position(saved_pos);
203            }
204            stream.set_position(saved_pos);
205            let span = stream
206                .peek()
207                .map(|(_, s)| s.clone())
208                .unwrap_or(Span { start: 0, end: 0 });
209            let expected = &[".f16", ".f32"];
210            let found = stream
211                .peek()
212                .map(|(t, _)| format!("{:?}", t))
213                .unwrap_or_else(|_| "<end of input>".to_string());
214            Err(crate::parser::unexpected_value(span, expected, found))
215        }
216    }
217
218    impl PtxParser for MmaSyncAlignedM8n8k4AlayoutBlayoutDtypeF16F16Ctype {
219        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
220            stream.expect_string("mma")?;
221            stream.expect_string(".sync")?;
222            let sync = ();
223            stream.expect_complete()?;
224            stream.expect_string(".aligned")?;
225            let aligned = ();
226            stream.expect_complete()?;
227            stream.expect_string(".m8n8k4")?;
228            let m8n8k4 = ();
229            stream.expect_complete()?;
230            let alayout = Alayout::parse(stream)?;
231            stream.expect_complete()?;
232            let blayout = Blayout::parse(stream)?;
233            stream.expect_complete()?;
234            let dtype = Dtype::parse(stream)?;
235            stream.expect_complete()?;
236            stream.expect_string(".f16")?;
237            let f16 = ();
238            stream.expect_complete()?;
239            stream.expect_string(".f16")?;
240            let f162 = ();
241            stream.expect_complete()?;
242            let ctype = Ctype::parse(stream)?;
243            stream.expect_complete()?;
244            let d = GeneralOperand::parse(stream)?;
245            stream.expect_complete()?;
246            stream.expect(&PtxToken::Comma)?;
247            let a = GeneralOperand::parse(stream)?;
248            stream.expect_complete()?;
249            stream.expect(&PtxToken::Comma)?;
250            let b = GeneralOperand::parse(stream)?;
251            stream.expect_complete()?;
252            stream.expect(&PtxToken::Comma)?;
253            let c = GeneralOperand::parse(stream)?;
254            stream.expect_complete()?;
255            stream.expect_complete()?;
256            stream.expect(&PtxToken::Semicolon)?;
257            Ok(MmaSyncAlignedM8n8k4AlayoutBlayoutDtypeF16F16Ctype {
258                sync,
259                aligned,
260                m8n8k4,
261                alayout,
262                blayout,
263                dtype,
264                f16,
265                f162,
266                ctype,
267                d,
268                a,
269                b,
270                c,
271            })
272        }
273    }
274
275    impl PtxParser for MmaSyncAlignedM16n8k8RowColDtypeF16F16Ctype {
276        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
277            stream.expect_string("mma")?;
278            stream.expect_string(".sync")?;
279            let sync = ();
280            stream.expect_complete()?;
281            stream.expect_string(".aligned")?;
282            let aligned = ();
283            stream.expect_complete()?;
284            stream.expect_string(".m16n8k8")?;
285            let m16n8k8 = ();
286            stream.expect_complete()?;
287            stream.expect_string(".row")?;
288            let row = ();
289            stream.expect_complete()?;
290            stream.expect_string(".col")?;
291            let col = ();
292            stream.expect_complete()?;
293            let dtype = Dtype::parse(stream)?;
294            stream.expect_complete()?;
295            stream.expect_string(".f16")?;
296            let f16 = ();
297            stream.expect_complete()?;
298            stream.expect_string(".f16")?;
299            let f162 = ();
300            stream.expect_complete()?;
301            let ctype = Ctype::parse(stream)?;
302            stream.expect_complete()?;
303            let d = GeneralOperand::parse(stream)?;
304            stream.expect_complete()?;
305            stream.expect(&PtxToken::Comma)?;
306            let a = GeneralOperand::parse(stream)?;
307            stream.expect_complete()?;
308            stream.expect(&PtxToken::Comma)?;
309            let b = GeneralOperand::parse(stream)?;
310            stream.expect_complete()?;
311            stream.expect(&PtxToken::Comma)?;
312            let c = GeneralOperand::parse(stream)?;
313            stream.expect_complete()?;
314            stream.expect_complete()?;
315            stream.expect(&PtxToken::Semicolon)?;
316            Ok(MmaSyncAlignedM16n8k8RowColDtypeF16F16Ctype {
317                sync,
318                aligned,
319                m16n8k8,
320                row,
321                col,
322                dtype,
323                f16,
324                f162,
325                ctype,
326                d,
327                a,
328                b,
329                c,
330            })
331        }
332    }
333
334    impl PtxParser for MmaSyncAlignedM16n8k16RowColDtypeF16F16Ctype {
335        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
336            stream.expect_string("mma")?;
337            stream.expect_string(".sync")?;
338            let sync = ();
339            stream.expect_complete()?;
340            stream.expect_string(".aligned")?;
341            let aligned = ();
342            stream.expect_complete()?;
343            stream.expect_string(".m16n8k16")?;
344            let m16n8k16 = ();
345            stream.expect_complete()?;
346            stream.expect_string(".row")?;
347            let row = ();
348            stream.expect_complete()?;
349            stream.expect_string(".col")?;
350            let col = ();
351            stream.expect_complete()?;
352            let dtype = Dtype::parse(stream)?;
353            stream.expect_complete()?;
354            stream.expect_string(".f16")?;
355            let f16 = ();
356            stream.expect_complete()?;
357            stream.expect_string(".f16")?;
358            let f162 = ();
359            stream.expect_complete()?;
360            let ctype = Ctype::parse(stream)?;
361            stream.expect_complete()?;
362            let d = GeneralOperand::parse(stream)?;
363            stream.expect_complete()?;
364            stream.expect(&PtxToken::Comma)?;
365            let a = GeneralOperand::parse(stream)?;
366            stream.expect_complete()?;
367            stream.expect(&PtxToken::Comma)?;
368            let b = GeneralOperand::parse(stream)?;
369            stream.expect_complete()?;
370            stream.expect(&PtxToken::Comma)?;
371            let c = GeneralOperand::parse(stream)?;
372            stream.expect_complete()?;
373            stream.expect_complete()?;
374            stream.expect(&PtxToken::Semicolon)?;
375            Ok(MmaSyncAlignedM16n8k16RowColDtypeF16F16Ctype {
376                sync,
377                aligned,
378                m16n8k16,
379                row,
380                col,
381                dtype,
382                f16,
383                f162,
384                ctype,
385                d,
386                a,
387                b,
388                c,
389            })
390        }
391    }
392}
393
394pub mod section_1 {
395    use super::*;
396    use crate::r#type::instruction::mma::section_1::*;
397
398    // ============================================================================
399    // Generated enum parsers
400    // ============================================================================
401
402    impl PtxParser for Atype {
403        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
404            // Try Bf16
405            {
406                let saved_pos = stream.position();
407                if stream.expect_string(".bf16").is_ok() {
408                    return Ok(Atype::Bf16);
409                }
410                stream.set_position(saved_pos);
411            }
412            let saved_pos = stream.position();
413            // Try Tf32
414            {
415                let saved_pos = stream.position();
416                if stream.expect_string(".tf32").is_ok() {
417                    return Ok(Atype::Tf32);
418                }
419                stream.set_position(saved_pos);
420            }
421            stream.set_position(saved_pos);
422            let span = stream
423                .peek()
424                .map(|(_, s)| s.clone())
425                .unwrap_or(Span { start: 0, end: 0 });
426            let expected = &[".bf16", ".tf32"];
427            let found = stream
428                .peek()
429                .map(|(t, _)| format!("{:?}", t))
430                .unwrap_or_else(|_| "<end of input>".to_string());
431            Err(crate::parser::unexpected_value(span, expected, found))
432        }
433    }
434
435    impl PtxParser for Btype {
436        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
437            // Try Bf16
438            {
439                let saved_pos = stream.position();
440                if stream.expect_string(".bf16").is_ok() {
441                    return Ok(Btype::Bf16);
442                }
443                stream.set_position(saved_pos);
444            }
445            let saved_pos = stream.position();
446            // Try Tf32
447            {
448                let saved_pos = stream.position();
449                if stream.expect_string(".tf32").is_ok() {
450                    return Ok(Btype::Tf32);
451                }
452                stream.set_position(saved_pos);
453            }
454            stream.set_position(saved_pos);
455            let span = stream
456                .peek()
457                .map(|(_, s)| s.clone())
458                .unwrap_or(Span { start: 0, end: 0 });
459            let expected = &[".bf16", ".tf32"];
460            let found = stream
461                .peek()
462                .map(|(t, _)| format!("{:?}", t))
463                .unwrap_or_else(|_| "<end of input>".to_string());
464            Err(crate::parser::unexpected_value(span, expected, found))
465        }
466    }
467
468    impl PtxParser for Ctype {
469        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
470            // Try F16
471            {
472                let saved_pos = stream.position();
473                if stream.expect_string(".f16").is_ok() {
474                    return Ok(Ctype::F16);
475                }
476                stream.set_position(saved_pos);
477            }
478            let saved_pos = stream.position();
479            // Try F32
480            {
481                let saved_pos = stream.position();
482                if stream.expect_string(".f32").is_ok() {
483                    return Ok(Ctype::F32);
484                }
485                stream.set_position(saved_pos);
486            }
487            stream.set_position(saved_pos);
488            let span = stream
489                .peek()
490                .map(|(_, s)| s.clone())
491                .unwrap_or(Span { start: 0, end: 0 });
492            let expected = &[".f16", ".f32"];
493            let found = stream
494                .peek()
495                .map(|(t, _)| format!("{:?}", t))
496                .unwrap_or_else(|_| "<end of input>".to_string());
497            Err(crate::parser::unexpected_value(span, expected, found))
498        }
499    }
500
501    impl PtxParser for Dtype {
502        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
503            // Try F16
504            {
505                let saved_pos = stream.position();
506                if stream.expect_string(".f16").is_ok() {
507                    return Ok(Dtype::F16);
508                }
509                stream.set_position(saved_pos);
510            }
511            let saved_pos = stream.position();
512            // Try F32
513            {
514                let saved_pos = stream.position();
515                if stream.expect_string(".f32").is_ok() {
516                    return Ok(Dtype::F32);
517                }
518                stream.set_position(saved_pos);
519            }
520            stream.set_position(saved_pos);
521            let span = stream
522                .peek()
523                .map(|(_, s)| s.clone())
524                .unwrap_or(Span { start: 0, end: 0 });
525            let expected = &[".f16", ".f32"];
526            let found = stream
527                .peek()
528                .map(|(t, _)| format!("{:?}", t))
529                .unwrap_or_else(|_| "<end of input>".to_string());
530            Err(crate::parser::unexpected_value(span, expected, found))
531        }
532    }
533
534    impl PtxParser for F8f6f4type {
535        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
536            // Try E4m3
537            {
538                let saved_pos = stream.position();
539                if stream.expect_string(".e4m3").is_ok() {
540                    return Ok(F8f6f4type::E4m3);
541                }
542                stream.set_position(saved_pos);
543            }
544            let saved_pos = stream.position();
545            // Try E5m2
546            {
547                let saved_pos = stream.position();
548                if stream.expect_string(".e5m2").is_ok() {
549                    return Ok(F8f6f4type::E5m2);
550                }
551                stream.set_position(saved_pos);
552            }
553            stream.set_position(saved_pos);
554            let saved_pos = stream.position();
555            // Try E3m2
556            {
557                let saved_pos = stream.position();
558                if stream.expect_string(".e3m2").is_ok() {
559                    return Ok(F8f6f4type::E3m2);
560                }
561                stream.set_position(saved_pos);
562            }
563            stream.set_position(saved_pos);
564            let saved_pos = stream.position();
565            // Try E2m3
566            {
567                let saved_pos = stream.position();
568                if stream.expect_string(".e2m3").is_ok() {
569                    return Ok(F8f6f4type::E2m3);
570                }
571                stream.set_position(saved_pos);
572            }
573            stream.set_position(saved_pos);
574            let saved_pos = stream.position();
575            // Try E2m1
576            {
577                let saved_pos = stream.position();
578                if stream.expect_string(".e2m1").is_ok() {
579                    return Ok(F8f6f4type::E2m1);
580                }
581                stream.set_position(saved_pos);
582            }
583            stream.set_position(saved_pos);
584            let span = stream
585                .peek()
586                .map(|(_, s)| s.clone())
587                .unwrap_or(Span { start: 0, end: 0 });
588            let expected = &[".e4m3", ".e5m2", ".e3m2", ".e2m3", ".e2m1"];
589            let found = stream
590                .peek()
591                .map(|(t, _)| format!("{:?}", t))
592                .unwrap_or_else(|_| "<end of input>".to_string());
593            Err(crate::parser::unexpected_value(span, expected, found))
594        }
595    }
596
597    impl PtxParser for F8type {
598        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
599            // Try E4m3
600            {
601                let saved_pos = stream.position();
602                if stream.expect_string(".e4m3").is_ok() {
603                    return Ok(F8type::E4m3);
604                }
605                stream.set_position(saved_pos);
606            }
607            let saved_pos = stream.position();
608            // Try E5m2
609            {
610                let saved_pos = stream.position();
611                if stream.expect_string(".e5m2").is_ok() {
612                    return Ok(F8type::E5m2);
613                }
614                stream.set_position(saved_pos);
615            }
616            stream.set_position(saved_pos);
617            let span = stream
618                .peek()
619                .map(|(_, s)| s.clone())
620                .unwrap_or(Span { start: 0, end: 0 });
621            let expected = &[".e4m3", ".e5m2"];
622            let found = stream
623                .peek()
624                .map(|(t, _)| format!("{:?}", t))
625                .unwrap_or_else(|_| "<end of input>".to_string());
626            Err(crate::parser::unexpected_value(span, expected, found))
627        }
628    }
629
630    impl PtxParser for Kind {
631        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
632            // Try KindF8f6f4
633            {
634                let saved_pos = stream.position();
635                if stream.expect_string(".kind::f8f6f4").is_ok() {
636                    return Ok(Kind::KindF8f6f4);
637                }
638                stream.set_position(saved_pos);
639            }
640            let span = stream
641                .peek()
642                .map(|(_, s)| s.clone())
643                .unwrap_or(Span { start: 0, end: 0 });
644            let expected = &[".kind::f8f6f4"];
645            let found = stream
646                .peek()
647                .map(|(t, _)| format!("{:?}", t))
648                .unwrap_or_else(|_| "<end of input>".to_string());
649            Err(crate::parser::unexpected_value(span, expected, found))
650        }
651    }
652
653    impl PtxParser for Shape {
654        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
655            // Try M16n8k16
656            {
657                let saved_pos = stream.position();
658                if stream.expect_string(".m16n8k16").is_ok() {
659                    return Ok(Shape::M16n8k16);
660                }
661                stream.set_position(saved_pos);
662            }
663            let saved_pos = stream.position();
664            // Try M16n8k32
665            {
666                let saved_pos = stream.position();
667                if stream.expect_string(".m16n8k32").is_ok() {
668                    return Ok(Shape::M16n8k32);
669                }
670                stream.set_position(saved_pos);
671            }
672            stream.set_position(saved_pos);
673            let span = stream
674                .peek()
675                .map(|(_, s)| s.clone())
676                .unwrap_or(Span { start: 0, end: 0 });
677            let expected = &[".m16n8k16", ".m16n8k32"];
678            let found = stream
679                .peek()
680                .map(|(t, _)| format!("{:?}", t))
681                .unwrap_or_else(|_| "<end of input>".to_string());
682            Err(crate::parser::unexpected_value(span, expected, found))
683        }
684    }
685
686    impl PtxParser for MmaSyncAlignedM16n8k4RowColF32Tf32Tf32F32 {
687        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
688            stream.expect_string("mma")?;
689            stream.expect_string(".sync")?;
690            let sync = ();
691            stream.expect_complete()?;
692            stream.expect_string(".aligned")?;
693            let aligned = ();
694            stream.expect_complete()?;
695            stream.expect_string(".m16n8k4")?;
696            let m16n8k4 = ();
697            stream.expect_complete()?;
698            stream.expect_string(".row")?;
699            let row = ();
700            stream.expect_complete()?;
701            stream.expect_string(".col")?;
702            let col = ();
703            stream.expect_complete()?;
704            stream.expect_string(".f32")?;
705            let f32 = ();
706            stream.expect_complete()?;
707            stream.expect_string(".tf32")?;
708            let tf32 = ();
709            stream.expect_complete()?;
710            stream.expect_string(".tf32")?;
711            let tf322 = ();
712            stream.expect_complete()?;
713            stream.expect_string(".f32")?;
714            let f322 = ();
715            stream.expect_complete()?;
716            let d = GeneralOperand::parse(stream)?;
717            stream.expect_complete()?;
718            stream.expect(&PtxToken::Comma)?;
719            let a = GeneralOperand::parse(stream)?;
720            stream.expect_complete()?;
721            stream.expect(&PtxToken::Comma)?;
722            let b = GeneralOperand::parse(stream)?;
723            stream.expect_complete()?;
724            stream.expect(&PtxToken::Comma)?;
725            let c = GeneralOperand::parse(stream)?;
726            stream.expect_complete()?;
727            stream.expect_complete()?;
728            stream.expect(&PtxToken::Semicolon)?;
729            Ok(MmaSyncAlignedM16n8k4RowColF32Tf32Tf32F32 {
730                sync,
731                aligned,
732                m16n8k4,
733                row,
734                col,
735                f32,
736                tf32,
737                tf322,
738                f322,
739                d,
740                a,
741                b,
742                c,
743            })
744        }
745    }
746
747    impl PtxParser for MmaSyncAlignedM16n8k8RowColF32AtypeBtypeF32 {
748        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
749            stream.expect_string("mma")?;
750            stream.expect_string(".sync")?;
751            let sync = ();
752            stream.expect_complete()?;
753            stream.expect_string(".aligned")?;
754            let aligned = ();
755            stream.expect_complete()?;
756            stream.expect_string(".m16n8k8")?;
757            let m16n8k8 = ();
758            stream.expect_complete()?;
759            stream.expect_string(".row")?;
760            let row = ();
761            stream.expect_complete()?;
762            stream.expect_string(".col")?;
763            let col = ();
764            stream.expect_complete()?;
765            stream.expect_string(".f32")?;
766            let f32 = ();
767            stream.expect_complete()?;
768            let atype = Atype::parse(stream)?;
769            stream.expect_complete()?;
770            let btype = Btype::parse(stream)?;
771            stream.expect_complete()?;
772            stream.expect_string(".f32")?;
773            let f322 = ();
774            stream.expect_complete()?;
775            let d = GeneralOperand::parse(stream)?;
776            stream.expect_complete()?;
777            stream.expect(&PtxToken::Comma)?;
778            let a = GeneralOperand::parse(stream)?;
779            stream.expect_complete()?;
780            stream.expect(&PtxToken::Comma)?;
781            let b = GeneralOperand::parse(stream)?;
782            stream.expect_complete()?;
783            stream.expect(&PtxToken::Comma)?;
784            let c = GeneralOperand::parse(stream)?;
785            stream.expect_complete()?;
786            stream.expect_complete()?;
787            stream.expect(&PtxToken::Semicolon)?;
788            Ok(MmaSyncAlignedM16n8k8RowColF32AtypeBtypeF32 {
789                sync,
790                aligned,
791                m16n8k8,
792                row,
793                col,
794                f32,
795                atype,
796                btype,
797                f322,
798                d,
799                a,
800                b,
801                c,
802            })
803        }
804    }
805
806    impl PtxParser for MmaSyncAlignedM16n8k16RowColF32Bf16Bf16F32 {
807        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
808            stream.expect_string("mma")?;
809            stream.expect_string(".sync")?;
810            let sync = ();
811            stream.expect_complete()?;
812            stream.expect_string(".aligned")?;
813            let aligned = ();
814            stream.expect_complete()?;
815            stream.expect_string(".m16n8k16")?;
816            let m16n8k16 = ();
817            stream.expect_complete()?;
818            stream.expect_string(".row")?;
819            let row = ();
820            stream.expect_complete()?;
821            stream.expect_string(".col")?;
822            let col = ();
823            stream.expect_complete()?;
824            stream.expect_string(".f32")?;
825            let f32 = ();
826            stream.expect_complete()?;
827            stream.expect_string(".bf16")?;
828            let bf16 = ();
829            stream.expect_complete()?;
830            stream.expect_string(".bf16")?;
831            let bf162 = ();
832            stream.expect_complete()?;
833            stream.expect_string(".f32")?;
834            let f322 = ();
835            stream.expect_complete()?;
836            let d = GeneralOperand::parse(stream)?;
837            stream.expect_complete()?;
838            stream.expect(&PtxToken::Comma)?;
839            let a = GeneralOperand::parse(stream)?;
840            stream.expect_complete()?;
841            stream.expect(&PtxToken::Comma)?;
842            let b = GeneralOperand::parse(stream)?;
843            stream.expect_complete()?;
844            stream.expect(&PtxToken::Comma)?;
845            let c = GeneralOperand::parse(stream)?;
846            stream.expect_complete()?;
847            stream.expect_complete()?;
848            stream.expect(&PtxToken::Semicolon)?;
849            Ok(MmaSyncAlignedM16n8k16RowColF32Bf16Bf16F32 {
850                sync,
851                aligned,
852                m16n8k16,
853                row,
854                col,
855                f32,
856                bf16,
857                bf162,
858                f322,
859                d,
860                a,
861                b,
862                c,
863            })
864        }
865    }
866
867    impl PtxParser for MmaSyncAlignedShapeRowColDtypeF8typeF8typeCtype {
868        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
869            stream.expect_string("mma")?;
870            stream.expect_string(".sync")?;
871            let sync = ();
872            stream.expect_complete()?;
873            stream.expect_string(".aligned")?;
874            let aligned = ();
875            stream.expect_complete()?;
876            let shape = Shape::parse(stream)?;
877            stream.expect_complete()?;
878            stream.expect_string(".row")?;
879            let row = ();
880            stream.expect_complete()?;
881            stream.expect_string(".col")?;
882            let col = ();
883            stream.expect_complete()?;
884            let dtype = Dtype::parse(stream)?;
885            stream.expect_complete()?;
886            let f8type = F8type::parse(stream)?;
887            stream.expect_complete()?;
888            let f8type1 = F8type::parse(stream)?;
889            stream.expect_complete()?;
890            let ctype = Ctype::parse(stream)?;
891            stream.expect_complete()?;
892            let d = GeneralOperand::parse(stream)?;
893            stream.expect_complete()?;
894            stream.expect(&PtxToken::Comma)?;
895            let a = GeneralOperand::parse(stream)?;
896            stream.expect_complete()?;
897            stream.expect(&PtxToken::Comma)?;
898            let b = GeneralOperand::parse(stream)?;
899            stream.expect_complete()?;
900            stream.expect(&PtxToken::Comma)?;
901            let c = GeneralOperand::parse(stream)?;
902            stream.expect_complete()?;
903            stream.expect_complete()?;
904            stream.expect(&PtxToken::Semicolon)?;
905            Ok(MmaSyncAlignedShapeRowColDtypeF8typeF8typeCtype {
906                sync,
907                aligned,
908                shape,
909                row,
910                col,
911                dtype,
912                f8type,
913                f8type1,
914                ctype,
915                d,
916                a,
917                b,
918                c,
919            })
920        }
921    }
922
923    impl PtxParser for MmaSyncAlignedM16n8k32RowColKindDtypeF8f6f4typeF8f6f4typeCtype {
924        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
925            stream.expect_string("mma")?;
926            stream.expect_string(".sync")?;
927            let sync = ();
928            stream.expect_complete()?;
929            stream.expect_string(".aligned")?;
930            let aligned = ();
931            stream.expect_complete()?;
932            stream.expect_string(".m16n8k32")?;
933            let m16n8k32 = ();
934            stream.expect_complete()?;
935            stream.expect_string(".row")?;
936            let row = ();
937            stream.expect_complete()?;
938            stream.expect_string(".col")?;
939            let col = ();
940            stream.expect_complete()?;
941            let kind = Kind::parse(stream)?;
942            stream.expect_complete()?;
943            let dtype = Dtype::parse(stream)?;
944            stream.expect_complete()?;
945            let f8f6f4type = F8f6f4type::parse(stream)?;
946            stream.expect_complete()?;
947            let f8f6f4type1 = F8f6f4type::parse(stream)?;
948            stream.expect_complete()?;
949            let ctype = Ctype::parse(stream)?;
950            stream.expect_complete()?;
951            let d = GeneralOperand::parse(stream)?;
952            stream.expect_complete()?;
953            stream.expect(&PtxToken::Comma)?;
954            let a = GeneralOperand::parse(stream)?;
955            stream.expect_complete()?;
956            stream.expect(&PtxToken::Comma)?;
957            let b = GeneralOperand::parse(stream)?;
958            stream.expect_complete()?;
959            stream.expect(&PtxToken::Comma)?;
960            let c = GeneralOperand::parse(stream)?;
961            stream.expect_complete()?;
962            stream.expect_complete()?;
963            stream.expect(&PtxToken::Semicolon)?;
964            Ok(
965                MmaSyncAlignedM16n8k32RowColKindDtypeF8f6f4typeF8f6f4typeCtype {
966                    sync,
967                    aligned,
968                    m16n8k32,
969                    row,
970                    col,
971                    kind,
972                    dtype,
973                    f8f6f4type,
974                    f8f6f4type1,
975                    ctype,
976                    d,
977                    a,
978                    b,
979                    c,
980                },
981            )
982        }
983    }
984}
985
986pub mod section_2 {
987    use super::*;
988    use crate::r#type::instruction::mma::section_2::*;
989
990    // ============================================================================
991    // Generated enum parsers
992    // ============================================================================
993
994    impl PtxParser for Kind {
995        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
996            // Try KindMxf4
997            {
998                let saved_pos = stream.position();
999                if stream.expect_string(".kind::mxf4").is_ok() {
1000                    return Ok(Kind::KindMxf4);
1001                }
1002                stream.set_position(saved_pos);
1003            }
1004            let span = stream
1005                .peek()
1006                .map(|(_, s)| s.clone())
1007                .unwrap_or(Span { start: 0, end: 0 });
1008            let expected = &[".kind::mxf4"];
1009            let found = stream
1010                .peek()
1011                .map(|(t, _)| format!("{:?}", t))
1012                .unwrap_or_else(|_| "<end of input>".to_string());
1013            Err(crate::parser::unexpected_value(span, expected, found))
1014        }
1015    }
1016
1017    impl PtxParser for ScaleVecSize {
1018        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1019            // Try ScaleVec2x
1020            {
1021                let saved_pos = stream.position();
1022                if stream.expect_string(".scale_vec::2X").is_ok() {
1023                    return Ok(ScaleVecSize::ScaleVec2x);
1024                }
1025                stream.set_position(saved_pos);
1026            }
1027            let span = stream
1028                .peek()
1029                .map(|(_, s)| s.clone())
1030                .unwrap_or(Span { start: 0, end: 0 });
1031            let expected = &[".scale_vec::2X"];
1032            let found = stream
1033                .peek()
1034                .map(|(t, _)| format!("{:?}", t))
1035                .unwrap_or_else(|_| "<end of input>".to_string());
1036            Err(crate::parser::unexpected_value(span, expected, found))
1037        }
1038    }
1039
1040    impl PtxParser for Stype {
1041        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1042            // Try Ue8m0
1043            {
1044                let saved_pos = stream.position();
1045                if stream.expect_string(".ue8m0").is_ok() {
1046                    return Ok(Stype::Ue8m0);
1047                }
1048                stream.set_position(saved_pos);
1049            }
1050            let span = stream
1051                .peek()
1052                .map(|(_, s)| s.clone())
1053                .unwrap_or(Span { start: 0, end: 0 });
1054            let expected = &[".ue8m0"];
1055            let found = stream
1056                .peek()
1057                .map(|(t, _)| format!("{:?}", t))
1058                .unwrap_or_else(|_| "<end of input>".to_string());
1059            Err(crate::parser::unexpected_value(span, expected, found))
1060        }
1061    }
1062
1063    impl PtxParser for MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype {
1064        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1065            stream.expect_string("mma")?;
1066            stream.expect_string(".sync")?;
1067            let sync = ();
1068            stream.expect_complete()?;
1069            stream.expect_string(".aligned")?;
1070            let aligned = ();
1071            stream.expect_complete()?;
1072            stream.expect_string(".m16n8k64")?;
1073            let m16n8k64 = ();
1074            stream.expect_complete()?;
1075            stream.expect_string(".row")?;
1076            let row = ();
1077            stream.expect_complete()?;
1078            stream.expect_string(".col")?;
1079            let col = ();
1080            stream.expect_complete()?;
1081            let kind = Kind::parse(stream)?;
1082            stream.expect_complete()?;
1083            stream.expect_string(".block_scale")?;
1084            let block_scale = ();
1085            stream.expect_complete()?;
1086            let saved_pos = stream.position();
1087            let scale_vec_size = match ScaleVecSize::parse(stream) {
1088                Ok(val) => Some(val),
1089                Err(_) => {
1090                    stream.set_position(saved_pos);
1091                    None
1092                }
1093            };
1094            stream.expect_complete()?;
1095            stream.expect_string(".f32")?;
1096            let f32 = ();
1097            stream.expect_complete()?;
1098            stream.expect_string(".e2m1")?;
1099            let e2m1 = ();
1100            stream.expect_complete()?;
1101            stream.expect_string(".e2m1")?;
1102            let e2m12 = ();
1103            stream.expect_complete()?;
1104            stream.expect_string(".f32")?;
1105            let f322 = ();
1106            stream.expect_complete()?;
1107            let stype = Stype::parse(stream)?;
1108            stream.expect_complete()?;
1109            let d = GeneralOperand::parse(stream)?;
1110            stream.expect_complete()?;
1111            stream.expect(&PtxToken::Comma)?;
1112            let a = GeneralOperand::parse(stream)?;
1113            stream.expect_complete()?;
1114            stream.expect(&PtxToken::Comma)?;
1115            let b = GeneralOperand::parse(stream)?;
1116            stream.expect_complete()?;
1117            stream.expect(&PtxToken::Comma)?;
1118            let c = GeneralOperand::parse(stream)?;
1119            stream.expect_complete()?;
1120            stream.expect(&PtxToken::Comma)?;
1121            let scale_a_data = GeneralOperand::parse(stream)?;
1122            stream.expect_complete()?;
1123            stream.expect(&PtxToken::Comma)?;
1124            let byte_id_a = VectorOperand::parse(stream)?;
1125            stream.expect_complete()?;
1126            stream.expect(&PtxToken::Comma)?;
1127            let scale_b_data = GeneralOperand::parse(stream)?;
1128            stream.expect_complete()?;
1129            stream.expect(&PtxToken::Comma)?;
1130            let byte_id_b = VectorOperand::parse(stream)?;
1131            stream.expect_complete()?;
1132            stream.expect_complete()?;
1133            stream.expect(&PtxToken::Semicolon)?;
1134            Ok(
1135                MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype {
1136                    sync,
1137                    aligned,
1138                    m16n8k64,
1139                    row,
1140                    col,
1141                    kind,
1142                    block_scale,
1143                    scale_vec_size,
1144                    f32,
1145                    e2m1,
1146                    e2m12,
1147                    f322,
1148                    stype,
1149                    d,
1150                    a,
1151                    b,
1152                    c,
1153                    scale_a_data,
1154                    byte_id_a,
1155                    scale_b_data,
1156                    byte_id_b,
1157                },
1158            )
1159        }
1160    }
1161}
1162
1163pub mod section_3 {
1164    use super::*;
1165    use crate::r#type::instruction::mma::section_3::*;
1166
1167    // ============================================================================
1168    // Generated enum parsers
1169    // ============================================================================
1170
1171    impl PtxParser for Kind {
1172        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1173            // Try KindMxf4nvf4
1174            {
1175                let saved_pos = stream.position();
1176                if stream.expect_string(".kind::mxf4nvf4").is_ok() {
1177                    return Ok(Kind::KindMxf4nvf4);
1178                }
1179                stream.set_position(saved_pos);
1180            }
1181            let span = stream
1182                .peek()
1183                .map(|(_, s)| s.clone())
1184                .unwrap_or(Span { start: 0, end: 0 });
1185            let expected = &[".kind::mxf4nvf4"];
1186            let found = stream
1187                .peek()
1188                .map(|(t, _)| format!("{:?}", t))
1189                .unwrap_or_else(|_| "<end of input>".to_string());
1190            Err(crate::parser::unexpected_value(span, expected, found))
1191        }
1192    }
1193
1194    impl PtxParser for ScaleVecSize {
1195        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1196            // Try ScaleVec2x
1197            {
1198                let saved_pos = stream.position();
1199                if stream.expect_string(".scale_vec::2X").is_ok() {
1200                    return Ok(ScaleVecSize::ScaleVec2x);
1201                }
1202                stream.set_position(saved_pos);
1203            }
1204            let saved_pos = stream.position();
1205            // Try ScaleVec4x
1206            {
1207                let saved_pos = stream.position();
1208                if stream.expect_string(".scale_vec::4X").is_ok() {
1209                    return Ok(ScaleVecSize::ScaleVec4x);
1210                }
1211                stream.set_position(saved_pos);
1212            }
1213            stream.set_position(saved_pos);
1214            let span = stream
1215                .peek()
1216                .map(|(_, s)| s.clone())
1217                .unwrap_or(Span { start: 0, end: 0 });
1218            let expected = &[".scale_vec::2X", ".scale_vec::4X"];
1219            let found = stream
1220                .peek()
1221                .map(|(t, _)| format!("{:?}", t))
1222                .unwrap_or_else(|_| "<end of input>".to_string());
1223            Err(crate::parser::unexpected_value(span, expected, found))
1224        }
1225    }
1226
1227    impl PtxParser for Stype {
1228        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1229            // Try Ue8m0
1230            {
1231                let saved_pos = stream.position();
1232                if stream.expect_string(".ue8m0").is_ok() {
1233                    return Ok(Stype::Ue8m0);
1234                }
1235                stream.set_position(saved_pos);
1236            }
1237            let saved_pos = stream.position();
1238            // Try Ue4m3
1239            {
1240                let saved_pos = stream.position();
1241                if stream.expect_string(".ue4m3").is_ok() {
1242                    return Ok(Stype::Ue4m3);
1243                }
1244                stream.set_position(saved_pos);
1245            }
1246            stream.set_position(saved_pos);
1247            let span = stream
1248                .peek()
1249                .map(|(_, s)| s.clone())
1250                .unwrap_or(Span { start: 0, end: 0 });
1251            let expected = &[".ue8m0", ".ue4m3"];
1252            let found = stream
1253                .peek()
1254                .map(|(t, _)| format!("{:?}", t))
1255                .unwrap_or_else(|_| "<end of input>".to_string());
1256            Err(crate::parser::unexpected_value(span, expected, found))
1257        }
1258    }
1259
1260    impl PtxParser for MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype1 {
1261        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1262            stream.expect_string("mma")?;
1263            stream.expect_string(".sync")?;
1264            let sync = ();
1265            stream.expect_complete()?;
1266            stream.expect_string(".aligned")?;
1267            let aligned = ();
1268            stream.expect_complete()?;
1269            stream.expect_string(".m16n8k64")?;
1270            let m16n8k64 = ();
1271            stream.expect_complete()?;
1272            stream.expect_string(".row")?;
1273            let row = ();
1274            stream.expect_complete()?;
1275            stream.expect_string(".col")?;
1276            let col = ();
1277            stream.expect_complete()?;
1278            let kind = Kind::parse(stream)?;
1279            stream.expect_complete()?;
1280            stream.expect_string(".block_scale")?;
1281            let block_scale = ();
1282            stream.expect_complete()?;
1283            let scale_vec_size = ScaleVecSize::parse(stream)?;
1284            stream.expect_complete()?;
1285            stream.expect_string(".f32")?;
1286            let f32 = ();
1287            stream.expect_complete()?;
1288            stream.expect_string(".e2m1")?;
1289            let e2m1 = ();
1290            stream.expect_complete()?;
1291            stream.expect_string(".e2m1")?;
1292            let e2m12 = ();
1293            stream.expect_complete()?;
1294            stream.expect_string(".f32")?;
1295            let f322 = ();
1296            stream.expect_complete()?;
1297            let stype = Stype::parse(stream)?;
1298            stream.expect_complete()?;
1299            let d = GeneralOperand::parse(stream)?;
1300            stream.expect_complete()?;
1301            stream.expect(&PtxToken::Comma)?;
1302            let a = GeneralOperand::parse(stream)?;
1303            stream.expect_complete()?;
1304            stream.expect(&PtxToken::Comma)?;
1305            let b = GeneralOperand::parse(stream)?;
1306            stream.expect_complete()?;
1307            stream.expect(&PtxToken::Comma)?;
1308            let c = GeneralOperand::parse(stream)?;
1309            stream.expect_complete()?;
1310            stream.expect(&PtxToken::Comma)?;
1311            let scale_a_data = GeneralOperand::parse(stream)?;
1312            stream.expect_complete()?;
1313            stream.expect(&PtxToken::Comma)?;
1314            let byte_id_a = VectorOperand::parse(stream)?;
1315            stream.expect_complete()?;
1316            stream.expect(&PtxToken::Comma)?;
1317            let scale_b_data = GeneralOperand::parse(stream)?;
1318            stream.expect_complete()?;
1319            stream.expect(&PtxToken::Comma)?;
1320            let byte_id_b = VectorOperand::parse(stream)?;
1321            stream.expect_complete()?;
1322            stream.expect_complete()?;
1323            stream.expect(&PtxToken::Semicolon)?;
1324            Ok(
1325                MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype1 {
1326                    sync,
1327                    aligned,
1328                    m16n8k64,
1329                    row,
1330                    col,
1331                    kind,
1332                    block_scale,
1333                    scale_vec_size,
1334                    f32,
1335                    e2m1,
1336                    e2m12,
1337                    f322,
1338                    stype,
1339                    d,
1340                    a,
1341                    b,
1342                    c,
1343                    scale_a_data,
1344                    byte_id_a,
1345                    scale_b_data,
1346                    byte_id_b,
1347                },
1348            )
1349        }
1350    }
1351}
1352
1353pub mod section_4 {
1354    use super::*;
1355    use crate::r#type::instruction::mma::section_4::*;
1356
1357    // ============================================================================
1358    // Generated enum parsers
1359    // ============================================================================
1360
1361    impl PtxParser for F8f6f4type {
1362        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1363            // Try E4m3
1364            {
1365                let saved_pos = stream.position();
1366                if stream.expect_string(".e4m3").is_ok() {
1367                    return Ok(F8f6f4type::E4m3);
1368                }
1369                stream.set_position(saved_pos);
1370            }
1371            let saved_pos = stream.position();
1372            // Try E5m2
1373            {
1374                let saved_pos = stream.position();
1375                if stream.expect_string(".e5m2").is_ok() {
1376                    return Ok(F8f6f4type::E5m2);
1377                }
1378                stream.set_position(saved_pos);
1379            }
1380            stream.set_position(saved_pos);
1381            let saved_pos = stream.position();
1382            // Try E3m2
1383            {
1384                let saved_pos = stream.position();
1385                if stream.expect_string(".e3m2").is_ok() {
1386                    return Ok(F8f6f4type::E3m2);
1387                }
1388                stream.set_position(saved_pos);
1389            }
1390            stream.set_position(saved_pos);
1391            let saved_pos = stream.position();
1392            // Try E2m3
1393            {
1394                let saved_pos = stream.position();
1395                if stream.expect_string(".e2m3").is_ok() {
1396                    return Ok(F8f6f4type::E2m3);
1397                }
1398                stream.set_position(saved_pos);
1399            }
1400            stream.set_position(saved_pos);
1401            let saved_pos = stream.position();
1402            // Try E2m1
1403            {
1404                let saved_pos = stream.position();
1405                if stream.expect_string(".e2m1").is_ok() {
1406                    return Ok(F8f6f4type::E2m1);
1407                }
1408                stream.set_position(saved_pos);
1409            }
1410            stream.set_position(saved_pos);
1411            let span = stream
1412                .peek()
1413                .map(|(_, s)| s.clone())
1414                .unwrap_or(Span { start: 0, end: 0 });
1415            let expected = &[".e4m3", ".e5m2", ".e3m2", ".e2m3", ".e2m1"];
1416            let found = stream
1417                .peek()
1418                .map(|(t, _)| format!("{:?}", t))
1419                .unwrap_or_else(|_| "<end of input>".to_string());
1420            Err(crate::parser::unexpected_value(span, expected, found))
1421        }
1422    }
1423
1424    impl PtxParser for Kind {
1425        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1426            // Try KindMxf8f6f4
1427            {
1428                let saved_pos = stream.position();
1429                if stream.expect_string(".kind::mxf8f6f4").is_ok() {
1430                    return Ok(Kind::KindMxf8f6f4);
1431                }
1432                stream.set_position(saved_pos);
1433            }
1434            let span = stream
1435                .peek()
1436                .map(|(_, s)| s.clone())
1437                .unwrap_or(Span { start: 0, end: 0 });
1438            let expected = &[".kind::mxf8f6f4"];
1439            let found = stream
1440                .peek()
1441                .map(|(t, _)| format!("{:?}", t))
1442                .unwrap_or_else(|_| "<end of input>".to_string());
1443            Err(crate::parser::unexpected_value(span, expected, found))
1444        }
1445    }
1446
1447    impl PtxParser for ScaleVecSize {
1448        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1449            // Try ScaleVec1x
1450            {
1451                let saved_pos = stream.position();
1452                if stream.expect_string(".scale_vec::1X").is_ok() {
1453                    return Ok(ScaleVecSize::ScaleVec1x);
1454                }
1455                stream.set_position(saved_pos);
1456            }
1457            let span = stream
1458                .peek()
1459                .map(|(_, s)| s.clone())
1460                .unwrap_or(Span { start: 0, end: 0 });
1461            let expected = &[".scale_vec::1X"];
1462            let found = stream
1463                .peek()
1464                .map(|(t, _)| format!("{:?}", t))
1465                .unwrap_or_else(|_| "<end of input>".to_string());
1466            Err(crate::parser::unexpected_value(span, expected, found))
1467        }
1468    }
1469
1470    impl PtxParser for Stype {
1471        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1472            // Try Ue8m0
1473            {
1474                let saved_pos = stream.position();
1475                if stream.expect_string(".ue8m0").is_ok() {
1476                    return Ok(Stype::Ue8m0);
1477                }
1478                stream.set_position(saved_pos);
1479            }
1480            let span = stream
1481                .peek()
1482                .map(|(_, s)| s.clone())
1483                .unwrap_or(Span { start: 0, end: 0 });
1484            let expected = &[".ue8m0"];
1485            let found = stream
1486                .peek()
1487                .map(|(t, _)| format!("{:?}", t))
1488                .unwrap_or_else(|_| "<end of input>".to_string());
1489            Err(crate::parser::unexpected_value(span, expected, found))
1490        }
1491    }
1492
1493    impl PtxParser
1494        for MmaSyncAlignedM16n8k32RowColKindBlockScaleScaleVecSizeF32F8f6f4typeF8f6f4typeF32Stype
1495    {
1496        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1497            stream.expect_string("mma")?;
1498            stream.expect_string(".sync")?;
1499            let sync = ();
1500            stream.expect_complete()?;
1501            stream.expect_string(".aligned")?;
1502            let aligned = ();
1503            stream.expect_complete()?;
1504            stream.expect_string(".m16n8k32")?;
1505            let m16n8k32 = ();
1506            stream.expect_complete()?;
1507            stream.expect_string(".row")?;
1508            let row = ();
1509            stream.expect_complete()?;
1510            stream.expect_string(".col")?;
1511            let col = ();
1512            stream.expect_complete()?;
1513            let kind = Kind::parse(stream)?;
1514            stream.expect_complete()?;
1515            stream.expect_string(".block_scale")?;
1516            let block_scale = ();
1517            stream.expect_complete()?;
1518            let saved_pos = stream.position();
1519            let scale_vec_size = match ScaleVecSize::parse(stream) {
1520                Ok(val) => Some(val),
1521                Err(_) => {
1522                    stream.set_position(saved_pos);
1523                    None
1524                }
1525            };
1526            stream.expect_complete()?;
1527            stream.expect_string(".f32")?;
1528            let f32 = ();
1529            stream.expect_complete()?;
1530            let f8f6f4type = F8f6f4type::parse(stream)?;
1531            stream.expect_complete()?;
1532            let f8f6f4type1 = F8f6f4type::parse(stream)?;
1533            stream.expect_complete()?;
1534            stream.expect_string(".f32")?;
1535            let f322 = ();
1536            stream.expect_complete()?;
1537            let stype = Stype::parse(stream)?;
1538            stream.expect_complete()?;
1539            let d = GeneralOperand::parse(stream)?;
1540            stream.expect_complete()?;
1541            stream.expect(&PtxToken::Comma)?;
1542            let a = GeneralOperand::parse(stream)?;
1543            stream.expect_complete()?;
1544            stream.expect(&PtxToken::Comma)?;
1545            let b = GeneralOperand::parse(stream)?;
1546            stream.expect_complete()?;
1547            stream.expect(&PtxToken::Comma)?;
1548            let c = GeneralOperand::parse(stream)?;
1549            stream.expect_complete()?;
1550            stream.expect(&PtxToken::Comma)?;
1551            let scale_a_data = GeneralOperand::parse(stream)?;
1552            stream.expect_complete()?;
1553            stream.expect(&PtxToken::Comma)?;
1554            let byte_id_a = VectorOperand::parse(stream)?;
1555            stream.expect_complete()?;
1556            stream.expect(&PtxToken::Comma)?;
1557            let scale_b_data = GeneralOperand::parse(stream)?;
1558            stream.expect_complete()?;
1559            stream.expect(&PtxToken::Comma)?;
1560            let byte_id_b = VectorOperand::parse(stream)?;
1561            stream.expect_complete()?;
1562            stream.expect_complete()?;
1563            stream.expect(&PtxToken::Semicolon)?;
1564            Ok(MmaSyncAlignedM16n8k32RowColKindBlockScaleScaleVecSizeF32F8f6f4typeF8f6f4typeF32Stype {
1565                sync,
1566                aligned,
1567                m16n8k32,
1568                row,
1569                col,
1570                kind,
1571                block_scale,
1572                scale_vec_size,
1573                f32,
1574                f8f6f4type,
1575                f8f6f4type1,
1576                f322,
1577                stype,
1578                d,
1579                a,
1580                b,
1581                c,
1582                scale_a_data,
1583                byte_id_a,
1584                scale_b_data,
1585                byte_id_b,
1586            })
1587        }
1588    }
1589}
1590
1591pub mod section_5 {
1592    use super::*;
1593    use crate::r#type::instruction::mma::section_5::*;
1594
1595    // ============================================================================
1596    // Generated enum parsers
1597    // ============================================================================
1598
1599    impl PtxParser for Shape {
1600        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1601            // Try M16n8k16
1602            {
1603                let saved_pos = stream.position();
1604                if stream.expect_string(".m16n8k16").is_ok() {
1605                    return Ok(Shape::M16n8k16);
1606                }
1607                stream.set_position(saved_pos);
1608            }
1609            let saved_pos = stream.position();
1610            // Try M16n8k4
1611            {
1612                let saved_pos = stream.position();
1613                if stream.expect_string(".m16n8k4").is_ok() {
1614                    return Ok(Shape::M16n8k4);
1615                }
1616                stream.set_position(saved_pos);
1617            }
1618            stream.set_position(saved_pos);
1619            let saved_pos = stream.position();
1620            // Try M16n8k8
1621            {
1622                let saved_pos = stream.position();
1623                if stream.expect_string(".m16n8k8").is_ok() {
1624                    return Ok(Shape::M16n8k8);
1625                }
1626                stream.set_position(saved_pos);
1627            }
1628            stream.set_position(saved_pos);
1629            let saved_pos = stream.position();
1630            // Try M8n84
1631            {
1632                let saved_pos = stream.position();
1633                if stream.expect_string(".m8n84").is_ok() {
1634                    return Ok(Shape::M8n84);
1635                }
1636                stream.set_position(saved_pos);
1637            }
1638            stream.set_position(saved_pos);
1639            let span = stream
1640                .peek()
1641                .map(|(_, s)| s.clone())
1642                .unwrap_or(Span { start: 0, end: 0 });
1643            let expected = &[".m16n8k16", ".m16n8k4", ".m16n8k8", ".m8n84"];
1644            let found = stream
1645                .peek()
1646                .map(|(t, _)| format!("{:?}", t))
1647                .unwrap_or_else(|_| "<end of input>".to_string());
1648            Err(crate::parser::unexpected_value(span, expected, found))
1649        }
1650    }
1651
1652    impl PtxParser for MmaSyncAlignedShapeRowColF64F64F64F64 {
1653        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1654            stream.expect_string("mma")?;
1655            stream.expect_string(".sync")?;
1656            let sync = ();
1657            stream.expect_complete()?;
1658            stream.expect_string(".aligned")?;
1659            let aligned = ();
1660            stream.expect_complete()?;
1661            let shape = Shape::parse(stream)?;
1662            stream.expect_complete()?;
1663            stream.expect_string(".row")?;
1664            let row = ();
1665            stream.expect_complete()?;
1666            stream.expect_string(".col")?;
1667            let col = ();
1668            stream.expect_complete()?;
1669            stream.expect_string(".f64")?;
1670            let f64 = ();
1671            stream.expect_complete()?;
1672            stream.expect_string(".f64")?;
1673            let f642 = ();
1674            stream.expect_complete()?;
1675            stream.expect_string(".f64")?;
1676            let f644 = ();
1677            stream.expect_complete()?;
1678            stream.expect_string(".f64")?;
1679            let f646 = ();
1680            stream.expect_complete()?;
1681            let d = GeneralOperand::parse(stream)?;
1682            stream.expect_complete()?;
1683            stream.expect(&PtxToken::Comma)?;
1684            let a = GeneralOperand::parse(stream)?;
1685            stream.expect_complete()?;
1686            stream.expect(&PtxToken::Comma)?;
1687            let b = GeneralOperand::parse(stream)?;
1688            stream.expect_complete()?;
1689            stream.expect(&PtxToken::Comma)?;
1690            let c = GeneralOperand::parse(stream)?;
1691            stream.expect_complete()?;
1692            stream.expect_complete()?;
1693            stream.expect(&PtxToken::Semicolon)?;
1694            Ok(MmaSyncAlignedShapeRowColF64F64F64F64 {
1695                sync,
1696                aligned,
1697                shape,
1698                row,
1699                col,
1700                f64,
1701                f642,
1702                f644,
1703                f646,
1704                d,
1705                a,
1706                b,
1707                c,
1708            })
1709        }
1710    }
1711}
1712
1713pub mod section_6 {
1714    use super::*;
1715    use crate::r#type::instruction::mma::section_6::*;
1716
1717    // ============================================================================
1718    // Generated enum parsers
1719    // ============================================================================
1720
1721    impl PtxParser for Atype {
1722        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1723            // Try U8
1724            {
1725                let saved_pos = stream.position();
1726                if stream.expect_string(".u8").is_ok() {
1727                    return Ok(Atype::U8);
1728                }
1729                stream.set_position(saved_pos);
1730            }
1731            let saved_pos = stream.position();
1732            // Try S8
1733            {
1734                let saved_pos = stream.position();
1735                if stream.expect_string(".s8").is_ok() {
1736                    return Ok(Atype::S8);
1737                }
1738                stream.set_position(saved_pos);
1739            }
1740            stream.set_position(saved_pos);
1741            let span = stream
1742                .peek()
1743                .map(|(_, s)| s.clone())
1744                .unwrap_or(Span { start: 0, end: 0 });
1745            let expected = &[".u8", ".s8"];
1746            let found = stream
1747                .peek()
1748                .map(|(t, _)| format!("{:?}", t))
1749                .unwrap_or_else(|_| "<end of input>".to_string());
1750            Err(crate::parser::unexpected_value(span, expected, found))
1751        }
1752    }
1753
1754    impl PtxParser for Btype {
1755        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1756            // Try U8
1757            {
1758                let saved_pos = stream.position();
1759                if stream.expect_string(".u8").is_ok() {
1760                    return Ok(Btype::U8);
1761                }
1762                stream.set_position(saved_pos);
1763            }
1764            let saved_pos = stream.position();
1765            // Try S8
1766            {
1767                let saved_pos = stream.position();
1768                if stream.expect_string(".s8").is_ok() {
1769                    return Ok(Btype::S8);
1770                }
1771                stream.set_position(saved_pos);
1772            }
1773            stream.set_position(saved_pos);
1774            let span = stream
1775                .peek()
1776                .map(|(_, s)| s.clone())
1777                .unwrap_or(Span { start: 0, end: 0 });
1778            let expected = &[".u8", ".s8"];
1779            let found = stream
1780                .peek()
1781                .map(|(t, _)| format!("{:?}", t))
1782                .unwrap_or_else(|_| "<end of input>".to_string());
1783            Err(crate::parser::unexpected_value(span, expected, found))
1784        }
1785    }
1786
1787    impl PtxParser for Shape {
1788        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1789            // Try M16n8k16
1790            {
1791                let saved_pos = stream.position();
1792                if stream.expect_string(".m16n8k16").is_ok() {
1793                    return Ok(Shape::M16n8k16);
1794                }
1795                stream.set_position(saved_pos);
1796            }
1797            let saved_pos = stream.position();
1798            // Try M16n8k32
1799            {
1800                let saved_pos = stream.position();
1801                if stream.expect_string(".m16n8k32").is_ok() {
1802                    return Ok(Shape::M16n8k32);
1803                }
1804                stream.set_position(saved_pos);
1805            }
1806            stream.set_position(saved_pos);
1807            let saved_pos = stream.position();
1808            // Try M8n8k16
1809            {
1810                let saved_pos = stream.position();
1811                if stream.expect_string(".m8n8k16").is_ok() {
1812                    return Ok(Shape::M8n8k16);
1813                }
1814                stream.set_position(saved_pos);
1815            }
1816            stream.set_position(saved_pos);
1817            let span = stream
1818                .peek()
1819                .map(|(_, s)| s.clone())
1820                .unwrap_or(Span { start: 0, end: 0 });
1821            let expected = &[".m16n8k16", ".m16n8k32", ".m8n8k16"];
1822            let found = stream
1823                .peek()
1824                .map(|(t, _)| format!("{:?}", t))
1825                .unwrap_or_else(|_| "<end of input>".to_string());
1826            Err(crate::parser::unexpected_value(span, expected, found))
1827        }
1828    }
1829
1830    impl PtxParser for MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS32 {
1831        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1832            stream.expect_string("mma")?;
1833            stream.expect_string(".sync")?;
1834            let sync = ();
1835            stream.expect_complete()?;
1836            stream.expect_string(".aligned")?;
1837            let aligned = ();
1838            stream.expect_complete()?;
1839            let shape = Shape::parse(stream)?;
1840            stream.expect_complete()?;
1841            stream.expect_string(".row")?;
1842            let row = ();
1843            stream.expect_complete()?;
1844            stream.expect_string(".col")?;
1845            let col = ();
1846            stream.expect_complete()?;
1847            let saved_pos = stream.position();
1848            let satfinite = stream.expect_string(".satfinite").is_ok();
1849            if !satfinite {
1850                stream.set_position(saved_pos);
1851            }
1852            stream.expect_complete()?;
1853            stream.expect_string(".s32")?;
1854            let s32 = ();
1855            stream.expect_complete()?;
1856            let atype = Atype::parse(stream)?;
1857            stream.expect_complete()?;
1858            let btype = Btype::parse(stream)?;
1859            stream.expect_complete()?;
1860            stream.expect_string(".s32")?;
1861            let s322 = ();
1862            stream.expect_complete()?;
1863            let d = GeneralOperand::parse(stream)?;
1864            stream.expect_complete()?;
1865            stream.expect(&PtxToken::Comma)?;
1866            let a = GeneralOperand::parse(stream)?;
1867            stream.expect_complete()?;
1868            stream.expect(&PtxToken::Comma)?;
1869            let b = GeneralOperand::parse(stream)?;
1870            stream.expect_complete()?;
1871            stream.expect(&PtxToken::Comma)?;
1872            let c = GeneralOperand::parse(stream)?;
1873            stream.expect_complete()?;
1874            stream.expect_complete()?;
1875            stream.expect(&PtxToken::Semicolon)?;
1876            Ok(MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS32 {
1877                sync,
1878                aligned,
1879                shape,
1880                row,
1881                col,
1882                satfinite,
1883                s32,
1884                atype,
1885                btype,
1886                s322,
1887                d,
1888                a,
1889                b,
1890                c,
1891            })
1892        }
1893    }
1894}
1895
1896pub mod section_7 {
1897    use super::*;
1898    use crate::r#type::instruction::mma::section_7::*;
1899
1900    // ============================================================================
1901    // Generated enum parsers
1902    // ============================================================================
1903
1904    impl PtxParser for Atype {
1905        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1906            // Try U4
1907            {
1908                let saved_pos = stream.position();
1909                if stream.expect_string(".u4").is_ok() {
1910                    return Ok(Atype::U4);
1911                }
1912                stream.set_position(saved_pos);
1913            }
1914            let saved_pos = stream.position();
1915            // Try S4
1916            {
1917                let saved_pos = stream.position();
1918                if stream.expect_string(".s4").is_ok() {
1919                    return Ok(Atype::S4);
1920                }
1921                stream.set_position(saved_pos);
1922            }
1923            stream.set_position(saved_pos);
1924            let span = stream
1925                .peek()
1926                .map(|(_, s)| s.clone())
1927                .unwrap_or(Span { start: 0, end: 0 });
1928            let expected = &[".u4", ".s4"];
1929            let found = stream
1930                .peek()
1931                .map(|(t, _)| format!("{:?}", t))
1932                .unwrap_or_else(|_| "<end of input>".to_string());
1933            Err(crate::parser::unexpected_value(span, expected, found))
1934        }
1935    }
1936
1937    impl PtxParser for Btype {
1938        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1939            // Try U4
1940            {
1941                let saved_pos = stream.position();
1942                if stream.expect_string(".u4").is_ok() {
1943                    return Ok(Btype::U4);
1944                }
1945                stream.set_position(saved_pos);
1946            }
1947            let saved_pos = stream.position();
1948            // Try S4
1949            {
1950                let saved_pos = stream.position();
1951                if stream.expect_string(".s4").is_ok() {
1952                    return Ok(Btype::S4);
1953                }
1954                stream.set_position(saved_pos);
1955            }
1956            stream.set_position(saved_pos);
1957            let span = stream
1958                .peek()
1959                .map(|(_, s)| s.clone())
1960                .unwrap_or(Span { start: 0, end: 0 });
1961            let expected = &[".u4", ".s4"];
1962            let found = stream
1963                .peek()
1964                .map(|(t, _)| format!("{:?}", t))
1965                .unwrap_or_else(|_| "<end of input>".to_string());
1966            Err(crate::parser::unexpected_value(span, expected, found))
1967        }
1968    }
1969
1970    impl PtxParser for Shape {
1971        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1972            // Try M16n8k32
1973            {
1974                let saved_pos = stream.position();
1975                if stream.expect_string(".m16n8k32").is_ok() {
1976                    return Ok(Shape::M16n8k32);
1977                }
1978                stream.set_position(saved_pos);
1979            }
1980            let saved_pos = stream.position();
1981            // Try M16n8k64
1982            {
1983                let saved_pos = stream.position();
1984                if stream.expect_string(".m16n8k64").is_ok() {
1985                    return Ok(Shape::M16n8k64);
1986                }
1987                stream.set_position(saved_pos);
1988            }
1989            stream.set_position(saved_pos);
1990            let saved_pos = stream.position();
1991            // Try M8n8k32
1992            {
1993                let saved_pos = stream.position();
1994                if stream.expect_string(".m8n8k32").is_ok() {
1995                    return Ok(Shape::M8n8k32);
1996                }
1997                stream.set_position(saved_pos);
1998            }
1999            stream.set_position(saved_pos);
2000            let span = stream
2001                .peek()
2002                .map(|(_, s)| s.clone())
2003                .unwrap_or(Span { start: 0, end: 0 });
2004            let expected = &[".m16n8k32", ".m16n8k64", ".m8n8k32"];
2005            let found = stream
2006                .peek()
2007                .map(|(t, _)| format!("{:?}", t))
2008                .unwrap_or_else(|_| "<end of input>".to_string());
2009            Err(crate::parser::unexpected_value(span, expected, found))
2010        }
2011    }
2012
2013    impl PtxParser for MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS321 {
2014        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2015            stream.expect_string("mma")?;
2016            stream.expect_string(".sync")?;
2017            let sync = ();
2018            stream.expect_complete()?;
2019            stream.expect_string(".aligned")?;
2020            let aligned = ();
2021            stream.expect_complete()?;
2022            let shape = Shape::parse(stream)?;
2023            stream.expect_complete()?;
2024            stream.expect_string(".row")?;
2025            let row = ();
2026            stream.expect_complete()?;
2027            stream.expect_string(".col")?;
2028            let col = ();
2029            stream.expect_complete()?;
2030            let saved_pos = stream.position();
2031            let satfinite = stream.expect_string(".satfinite").is_ok();
2032            if !satfinite {
2033                stream.set_position(saved_pos);
2034            }
2035            stream.expect_complete()?;
2036            stream.expect_string(".s32")?;
2037            let s32 = ();
2038            stream.expect_complete()?;
2039            let atype = Atype::parse(stream)?;
2040            stream.expect_complete()?;
2041            let btype = Btype::parse(stream)?;
2042            stream.expect_complete()?;
2043            stream.expect_string(".s32")?;
2044            let s322 = ();
2045            stream.expect_complete()?;
2046            let d = GeneralOperand::parse(stream)?;
2047            stream.expect_complete()?;
2048            stream.expect(&PtxToken::Comma)?;
2049            let a = GeneralOperand::parse(stream)?;
2050            stream.expect_complete()?;
2051            stream.expect(&PtxToken::Comma)?;
2052            let b = GeneralOperand::parse(stream)?;
2053            stream.expect_complete()?;
2054            stream.expect(&PtxToken::Comma)?;
2055            let c = GeneralOperand::parse(stream)?;
2056            stream.expect_complete()?;
2057            stream.expect_complete()?;
2058            stream.expect(&PtxToken::Semicolon)?;
2059            Ok(MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS321 {
2060                sync,
2061                aligned,
2062                shape,
2063                row,
2064                col,
2065                satfinite,
2066                s32,
2067                atype,
2068                btype,
2069                s322,
2070                d,
2071                a,
2072                b,
2073                c,
2074            })
2075        }
2076    }
2077}
2078
2079pub mod section_8 {
2080    use super::*;
2081    use crate::r#type::instruction::mma::section_8::*;
2082
2083    // ============================================================================
2084    // Generated enum parsers
2085    // ============================================================================
2086
2087    impl PtxParser for Bitop {
2088        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2089            // Try Xor
2090            {
2091                let saved_pos = stream.position();
2092                if stream.expect_string(".xor").is_ok() {
2093                    return Ok(Bitop::Xor);
2094                }
2095                stream.set_position(saved_pos);
2096            }
2097            let saved_pos = stream.position();
2098            // Try And
2099            {
2100                let saved_pos = stream.position();
2101                if stream.expect_string(".and").is_ok() {
2102                    return Ok(Bitop::And);
2103                }
2104                stream.set_position(saved_pos);
2105            }
2106            stream.set_position(saved_pos);
2107            let span = stream
2108                .peek()
2109                .map(|(_, s)| s.clone())
2110                .unwrap_or(Span { start: 0, end: 0 });
2111            let expected = &[".xor", ".and"];
2112            let found = stream
2113                .peek()
2114                .map(|(t, _)| format!("{:?}", t))
2115                .unwrap_or_else(|_| "<end of input>".to_string());
2116            Err(crate::parser::unexpected_value(span, expected, found))
2117        }
2118    }
2119
2120    impl PtxParser for Shape {
2121        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2122            // Try M16n8k128
2123            {
2124                let saved_pos = stream.position();
2125                if stream.expect_string(".m16n8k128").is_ok() {
2126                    return Ok(Shape::M16n8k128);
2127                }
2128                stream.set_position(saved_pos);
2129            }
2130            let saved_pos = stream.position();
2131            // Try M16n8k256
2132            {
2133                let saved_pos = stream.position();
2134                if stream.expect_string(".m16n8k256").is_ok() {
2135                    return Ok(Shape::M16n8k256);
2136                }
2137                stream.set_position(saved_pos);
2138            }
2139            stream.set_position(saved_pos);
2140            let saved_pos = stream.position();
2141            // Try M8n8k128
2142            {
2143                let saved_pos = stream.position();
2144                if stream.expect_string(".m8n8k128").is_ok() {
2145                    return Ok(Shape::M8n8k128);
2146                }
2147                stream.set_position(saved_pos);
2148            }
2149            stream.set_position(saved_pos);
2150            let span = stream
2151                .peek()
2152                .map(|(_, s)| s.clone())
2153                .unwrap_or(Span { start: 0, end: 0 });
2154            let expected = &[".m16n8k128", ".m16n8k256", ".m8n8k128"];
2155            let found = stream
2156                .peek()
2157                .map(|(t, _)| format!("{:?}", t))
2158                .unwrap_or_else(|_| "<end of input>".to_string());
2159            Err(crate::parser::unexpected_value(span, expected, found))
2160        }
2161    }
2162
2163    impl PtxParser for MmaSyncAlignedShapeRowColS32B1B1S32BitopPopc {
2164        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2165            stream.expect_string("mma")?;
2166            stream.expect_string(".sync")?;
2167            let sync = ();
2168            stream.expect_complete()?;
2169            stream.expect_string(".aligned")?;
2170            let aligned = ();
2171            stream.expect_complete()?;
2172            let shape = Shape::parse(stream)?;
2173            stream.expect_complete()?;
2174            stream.expect_string(".row")?;
2175            let row = ();
2176            stream.expect_complete()?;
2177            stream.expect_string(".col")?;
2178            let col = ();
2179            stream.expect_complete()?;
2180            stream.expect_string(".s32")?;
2181            let s32 = ();
2182            stream.expect_complete()?;
2183            stream.expect_string(".b1")?;
2184            let b1 = ();
2185            stream.expect_complete()?;
2186            stream.expect_string(".b1")?;
2187            let b12 = ();
2188            stream.expect_complete()?;
2189            stream.expect_string(".s32")?;
2190            let s322 = ();
2191            stream.expect_complete()?;
2192            let bitop = Bitop::parse(stream)?;
2193            stream.expect_complete()?;
2194            stream.expect_string(".popc")?;
2195            let popc = ();
2196            stream.expect_complete()?;
2197            let d = GeneralOperand::parse(stream)?;
2198            stream.expect_complete()?;
2199            stream.expect(&PtxToken::Comma)?;
2200            let a = GeneralOperand::parse(stream)?;
2201            stream.expect_complete()?;
2202            stream.expect(&PtxToken::Comma)?;
2203            let b = GeneralOperand::parse(stream)?;
2204            stream.expect_complete()?;
2205            stream.expect(&PtxToken::Comma)?;
2206            let c = GeneralOperand::parse(stream)?;
2207            stream.expect_complete()?;
2208            stream.expect_complete()?;
2209            stream.expect(&PtxToken::Semicolon)?;
2210            Ok(MmaSyncAlignedShapeRowColS32B1B1S32BitopPopc {
2211                sync,
2212                aligned,
2213                shape,
2214                row,
2215                col,
2216                s32,
2217                b1,
2218                b12,
2219                s322,
2220                bitop,
2221                popc,
2222                d,
2223                a,
2224                b,
2225                c,
2226            })
2227        }
2228    }
2229}
ptx_parser/parser/instruction/mma.rs

ptx_parser/parser/instruction/
mma.rs
