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

ptx_parser/parser/instruction/
mma.rs
