ptx_parser/parser/instruction/
wgmma_mma_async.rs

1//! Original PTX specification:
2//!
3//! // Half precision floating point type:
4//! wgmma.mma_async.sync.aligned.shape.dtype.f16.f16  d, a-desc, b-desc, scale-d, imm-scale-a, imm-scale-b, imm-trans-a, imm-trans-b;
5//! wgmma.mma_async.sync.aligned.shape.dtype.f16.f16  d, a, b-desc, scale-d, imm-scale-a, imm-scale-b, imm-trans-b;
6//! .shape   = {.m64n8k16, .m64n16k16, .m64n24k16, .m64n32k16,
7//! .m64n40k16, .m64n48k16, .m64n56k16, .m64n64k16,
8//! .m64n72k16, .m64n80k16, .m64n88k16, .m64n96k16,
9//! .m64n104k16, .m64n112k16, .m64n120k16, .m64n128k16,
10//! .m64n136k16, .m64n144k16, .m64n152k16, .m64n160k16,
11//! .m64n168k16, .m64n176k16, .m64n184k16, .m64n192k16,
12//! .m64n200k16, .m64n208k16, .m64n216k16, .m64n224k16,
13//! .m64n232k16, .m64n240k16, .m64n248k16, .m64n256k16};
14//! .dtype   = {.f16, .f32};
15//! ------------------------------------------------------------------
16//! // Alternate floating point type :
17//! // .bf16 floating point type:
18//! wgmma.mma_async.sync.aligned.shape.dtype.bf16.bf16  d, a-desc, b-desc, scale-d, imm-scale-a, imm-scale-b, imm-trans-a, imm-trans-b;
19//! wgmma.mma_async.sync.aligned.shape.dtype.bf16.bf16  d, a, b-desc, scale-d, imm-scale-a, imm-scale-b, imm-trans-b;
20//! .shape   = {.m64n8k16, .m64n16k16, .m64n24k16, .m64n32k16,
21//! .m64n40k16, .m64n48k16, .m64n56k16, .m64n64k16,
22//! .m64n72k16, .m64n80k16, .m64n88k16, .m64n96k16,
23//! .m64n104k16, .m64n112k16, .m64n120k16, .m64n128k16,
24//! .m64n136k16, .m64n144k16, .m64n152k16, .m64n160k16,
25//! .m64n168k16, .m64n176k16, .m64n184k16, .m64n192k16,
26//! .m64n200k16, .m64n208k16, .m64n216k16, .m64n224k16,
27//! .m64n232k16, .m64n240k16, .m64n248k16, .m64n256k16};
28//! .dtype  = {.f32};
29//! ------------------------------------------------------------------
30//! // .tf32 floating point type:
31//! wgmma.mma_async.sync.aligned.shape.dtype.tf32.tf32  d, a-desc, b-desc, scale-d, imm-scale-a, imm-scale-b;
32//! wgmma.mma_async.sync.aligned.shape.dtype.tf32.tf32  d, a, b-desc, scale-d, imm-scale-a, imm-scale-b;
33//! .shape   = {.m64n8k8, .m64n16k8, .m64n24k8, .m64n32k8,
34//! .m64n40k8, .m64n48k8, .m64n56k8, .m64n64k8,
35//! .m64n72k8, .m64n80k8, .m64n88k8, .m64n96k8,
36//! .m64n104k8, .m64n112k8, .m64n120k8, .m64n128k8,
37//! .m64n136k8, .m64n144k8, .m64n152k8, .m64n160k8,
38//! .m64n168k8, .m64n176k8, .m64n184k8, .m64n192k8,
39//! .m64n200k8, .m64n208k8, .m64n216k8, .m64n224k8,
40//! .m64n232k8, .m64n240k8, .m64n248k8, .m64n256k8};
41//! .dtype  = {.f32};
42//! ------------------------------------------------------------------
43//! // FP8 floating point type
44//! wgmma.mma_async.sync.aligned.shape.dtype.atype.btype  d, a-desc, b-desc, scale-d, imm-scale-a, imm-scale-b;
45//! wgmma.mma_async.sync.aligned.shape.dtype.atype.btype  d, a, b-desc, scale-d, imm-scale-a, imm-scale-b;
46//! .shape   = {.m64n8k32, .m64n16k32, .m64n24k32, .m64n32k32,
47//! .m64n40k32, .m64n48k32, .m64n56k32, .m64n64k32,
48//! .m64n72k32, .m64n80k32, .m64n88k32, .m64n96k32,
49//! .m64n104k32, .m64n112k32, .m64n120k32, .m64n128k32,
50//! .m64n136k32, .m64n144k32, .m64n152k32, .m64n160k32,
51//! .m64n168k32, .m64n176k32, .m64n184k32, .m64n192k32,
52//! .m64n200k32, .m64n208k32, .m64n216k32, .m64n224k32,
53//! .m64n232k32, .m64n240k32, .m64n248k32, .m64n256k32};
54//! .atype  = {.e4m3, .e5m2};
55//! .btype  = {.e4m3, .e5m2};
56//! .dtype  = {.f16, .f32};
57//! ------------------------------------------------------------------
58//! // Integer type:
59//! wgmma.mma_async.sync.aligned.shape{.satfinite}.s32.atype.btype  d, a-desc, b-desc, scale-d;
60//! wgmma.mma_async.sync.aligned.shape{.satfinite}.s32.atype.btype  d, a, b-desc, scale-d;
61//! .shape   = {.m64n8k32, .m64n16k32, .m64n24k32, .m64n32k32,
62//! .m64n48k32, .m64n64k32, .m64n80k32, .m64n96k32,
63//! .m64n112k32, .m64n128k32, .m64n144k32, .m64n160k32,
64//! .m64n176k32, .m64n192k32, .m64n208k32, .m64n224k32};
65//! .atype  = {.s8, .u8};
66//! .btype  = {.s8, .u8};
67//! ------------------------------------------------------------------
68//! // Single bit:
69//! wgmma.mma_async.sync.aligned.shape.s32.b1.b1.op.popc  d, a-desc, b-desc, scale-d;
70//! wgmma.mma_async.sync.aligned.shape.s32.b1.b1.op.popc  d, a, b-desc, scale-d;
71//! .shape   = {.m64n8k256, .m64n16k256, .m64n24k256, .m64n32k256,
72//! .m64n48k256, .m64n64k256, .m64n80k256, .m64n96k256,
73//! .m64n112k256, .m64n128k256, .m64n144k256, .m64n160k256,
74//! .m64n176k256, .m64n192k256, .m64n208k256, .m64n224k256,
75//! .m64n240k256, .m64n256k256};
76//! .op  = {.and};
77
78#![allow(unused)]
79
80use crate::lexer::PtxToken;
81use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
82use crate::r#type::common::*;
83
84pub mod section_0 {
85    use super::*;
86    use crate::r#type::instruction::wgmma_mma_async::section_0::*;
87
88    // ============================================================================
89    // Generated enum parsers
90    // ============================================================================
91
92    impl PtxParser for Dtype {
93        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
94            // Try F16
95            {
96                let saved_pos = stream.position();
97                if stream.expect_string(".f16").is_ok() {
98                    return Ok(Dtype::F16);
99                }
100                stream.set_position(saved_pos);
101            }
102            let saved_pos = stream.position();
103            // Try F32
104            {
105                let saved_pos = stream.position();
106                if stream.expect_string(".f32").is_ok() {
107                    return Ok(Dtype::F32);
108                }
109                stream.set_position(saved_pos);
110            }
111            stream.set_position(saved_pos);
112            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
113            let expected = &[".f16", ".f32"];
114            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
115            Err(crate::parser::unexpected_value(span, expected, found))
116        }
117    }
118
119    impl PtxParser for Shape {
120        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
121            // Try M64n104k16
122            {
123                let saved_pos = stream.position();
124                if stream.expect_string(".m64n104k16").is_ok() {
125                    return Ok(Shape::M64n104k16);
126                }
127                stream.set_position(saved_pos);
128            }
129            let saved_pos = stream.position();
130            // Try M64n112k16
131            {
132                let saved_pos = stream.position();
133                if stream.expect_string(".m64n112k16").is_ok() {
134                    return Ok(Shape::M64n112k16);
135                }
136                stream.set_position(saved_pos);
137            }
138            stream.set_position(saved_pos);
139            let saved_pos = stream.position();
140            // Try M64n120k16
141            {
142                let saved_pos = stream.position();
143                if stream.expect_string(".m64n120k16").is_ok() {
144                    return Ok(Shape::M64n120k16);
145                }
146                stream.set_position(saved_pos);
147            }
148            stream.set_position(saved_pos);
149            let saved_pos = stream.position();
150            // Try M64n128k16
151            {
152                let saved_pos = stream.position();
153                if stream.expect_string(".m64n128k16").is_ok() {
154                    return Ok(Shape::M64n128k16);
155                }
156                stream.set_position(saved_pos);
157            }
158            stream.set_position(saved_pos);
159            let saved_pos = stream.position();
160            // Try M64n136k16
161            {
162                let saved_pos = stream.position();
163                if stream.expect_string(".m64n136k16").is_ok() {
164                    return Ok(Shape::M64n136k16);
165                }
166                stream.set_position(saved_pos);
167            }
168            stream.set_position(saved_pos);
169            let saved_pos = stream.position();
170            // Try M64n144k16
171            {
172                let saved_pos = stream.position();
173                if stream.expect_string(".m64n144k16").is_ok() {
174                    return Ok(Shape::M64n144k16);
175                }
176                stream.set_position(saved_pos);
177            }
178            stream.set_position(saved_pos);
179            let saved_pos = stream.position();
180            // Try M64n152k16
181            {
182                let saved_pos = stream.position();
183                if stream.expect_string(".m64n152k16").is_ok() {
184                    return Ok(Shape::M64n152k16);
185                }
186                stream.set_position(saved_pos);
187            }
188            stream.set_position(saved_pos);
189            let saved_pos = stream.position();
190            // Try M64n160k16
191            {
192                let saved_pos = stream.position();
193                if stream.expect_string(".m64n160k16").is_ok() {
194                    return Ok(Shape::M64n160k16);
195                }
196                stream.set_position(saved_pos);
197            }
198            stream.set_position(saved_pos);
199            let saved_pos = stream.position();
200            // Try M64n168k16
201            {
202                let saved_pos = stream.position();
203                if stream.expect_string(".m64n168k16").is_ok() {
204                    return Ok(Shape::M64n168k16);
205                }
206                stream.set_position(saved_pos);
207            }
208            stream.set_position(saved_pos);
209            let saved_pos = stream.position();
210            // Try M64n176k16
211            {
212                let saved_pos = stream.position();
213                if stream.expect_string(".m64n176k16").is_ok() {
214                    return Ok(Shape::M64n176k16);
215                }
216                stream.set_position(saved_pos);
217            }
218            stream.set_position(saved_pos);
219            let saved_pos = stream.position();
220            // Try M64n184k16
221            {
222                let saved_pos = stream.position();
223                if stream.expect_string(".m64n184k16").is_ok() {
224                    return Ok(Shape::M64n184k16);
225                }
226                stream.set_position(saved_pos);
227            }
228            stream.set_position(saved_pos);
229            let saved_pos = stream.position();
230            // Try M64n192k16
231            {
232                let saved_pos = stream.position();
233                if stream.expect_string(".m64n192k16").is_ok() {
234                    return Ok(Shape::M64n192k16);
235                }
236                stream.set_position(saved_pos);
237            }
238            stream.set_position(saved_pos);
239            let saved_pos = stream.position();
240            // Try M64n200k16
241            {
242                let saved_pos = stream.position();
243                if stream.expect_string(".m64n200k16").is_ok() {
244                    return Ok(Shape::M64n200k16);
245                }
246                stream.set_position(saved_pos);
247            }
248            stream.set_position(saved_pos);
249            let saved_pos = stream.position();
250            // Try M64n208k16
251            {
252                let saved_pos = stream.position();
253                if stream.expect_string(".m64n208k16").is_ok() {
254                    return Ok(Shape::M64n208k16);
255                }
256                stream.set_position(saved_pos);
257            }
258            stream.set_position(saved_pos);
259            let saved_pos = stream.position();
260            // Try M64n216k16
261            {
262                let saved_pos = stream.position();
263                if stream.expect_string(".m64n216k16").is_ok() {
264                    return Ok(Shape::M64n216k16);
265                }
266                stream.set_position(saved_pos);
267            }
268            stream.set_position(saved_pos);
269            let saved_pos = stream.position();
270            // Try M64n224k16
271            {
272                let saved_pos = stream.position();
273                if stream.expect_string(".m64n224k16").is_ok() {
274                    return Ok(Shape::M64n224k16);
275                }
276                stream.set_position(saved_pos);
277            }
278            stream.set_position(saved_pos);
279            let saved_pos = stream.position();
280            // Try M64n232k16
281            {
282                let saved_pos = stream.position();
283                if stream.expect_string(".m64n232k16").is_ok() {
284                    return Ok(Shape::M64n232k16);
285                }
286                stream.set_position(saved_pos);
287            }
288            stream.set_position(saved_pos);
289            let saved_pos = stream.position();
290            // Try M64n240k16
291            {
292                let saved_pos = stream.position();
293                if stream.expect_string(".m64n240k16").is_ok() {
294                    return Ok(Shape::M64n240k16);
295                }
296                stream.set_position(saved_pos);
297            }
298            stream.set_position(saved_pos);
299            let saved_pos = stream.position();
300            // Try M64n248k16
301            {
302                let saved_pos = stream.position();
303                if stream.expect_string(".m64n248k16").is_ok() {
304                    return Ok(Shape::M64n248k16);
305                }
306                stream.set_position(saved_pos);
307            }
308            stream.set_position(saved_pos);
309            let saved_pos = stream.position();
310            // Try M64n256k16
311            {
312                let saved_pos = stream.position();
313                if stream.expect_string(".m64n256k16").is_ok() {
314                    return Ok(Shape::M64n256k16);
315                }
316                stream.set_position(saved_pos);
317            }
318            stream.set_position(saved_pos);
319            let saved_pos = stream.position();
320            // Try M64n16k16
321            {
322                let saved_pos = stream.position();
323                if stream.expect_string(".m64n16k16").is_ok() {
324                    return Ok(Shape::M64n16k16);
325                }
326                stream.set_position(saved_pos);
327            }
328            stream.set_position(saved_pos);
329            let saved_pos = stream.position();
330            // Try M64n24k16
331            {
332                let saved_pos = stream.position();
333                if stream.expect_string(".m64n24k16").is_ok() {
334                    return Ok(Shape::M64n24k16);
335                }
336                stream.set_position(saved_pos);
337            }
338            stream.set_position(saved_pos);
339            let saved_pos = stream.position();
340            // Try M64n32k16
341            {
342                let saved_pos = stream.position();
343                if stream.expect_string(".m64n32k16").is_ok() {
344                    return Ok(Shape::M64n32k16);
345                }
346                stream.set_position(saved_pos);
347            }
348            stream.set_position(saved_pos);
349            let saved_pos = stream.position();
350            // Try M64n40k16
351            {
352                let saved_pos = stream.position();
353                if stream.expect_string(".m64n40k16").is_ok() {
354                    return Ok(Shape::M64n40k16);
355                }
356                stream.set_position(saved_pos);
357            }
358            stream.set_position(saved_pos);
359            let saved_pos = stream.position();
360            // Try M64n48k16
361            {
362                let saved_pos = stream.position();
363                if stream.expect_string(".m64n48k16").is_ok() {
364                    return Ok(Shape::M64n48k16);
365                }
366                stream.set_position(saved_pos);
367            }
368            stream.set_position(saved_pos);
369            let saved_pos = stream.position();
370            // Try M64n56k16
371            {
372                let saved_pos = stream.position();
373                if stream.expect_string(".m64n56k16").is_ok() {
374                    return Ok(Shape::M64n56k16);
375                }
376                stream.set_position(saved_pos);
377            }
378            stream.set_position(saved_pos);
379            let saved_pos = stream.position();
380            // Try M64n64k16
381            {
382                let saved_pos = stream.position();
383                if stream.expect_string(".m64n64k16").is_ok() {
384                    return Ok(Shape::M64n64k16);
385                }
386                stream.set_position(saved_pos);
387            }
388            stream.set_position(saved_pos);
389            let saved_pos = stream.position();
390            // Try M64n72k16
391            {
392                let saved_pos = stream.position();
393                if stream.expect_string(".m64n72k16").is_ok() {
394                    return Ok(Shape::M64n72k16);
395                }
396                stream.set_position(saved_pos);
397            }
398            stream.set_position(saved_pos);
399            let saved_pos = stream.position();
400            // Try M64n80k16
401            {
402                let saved_pos = stream.position();
403                if stream.expect_string(".m64n80k16").is_ok() {
404                    return Ok(Shape::M64n80k16);
405                }
406                stream.set_position(saved_pos);
407            }
408            stream.set_position(saved_pos);
409            let saved_pos = stream.position();
410            // Try M64n88k16
411            {
412                let saved_pos = stream.position();
413                if stream.expect_string(".m64n88k16").is_ok() {
414                    return Ok(Shape::M64n88k16);
415                }
416                stream.set_position(saved_pos);
417            }
418            stream.set_position(saved_pos);
419            let saved_pos = stream.position();
420            // Try M64n96k16
421            {
422                let saved_pos = stream.position();
423                if stream.expect_string(".m64n96k16").is_ok() {
424                    return Ok(Shape::M64n96k16);
425                }
426                stream.set_position(saved_pos);
427            }
428            stream.set_position(saved_pos);
429            let saved_pos = stream.position();
430            // Try M64n8k16
431            {
432                let saved_pos = stream.position();
433                if stream.expect_string(".m64n8k16").is_ok() {
434                    return Ok(Shape::M64n8k16);
435                }
436                stream.set_position(saved_pos);
437            }
438            stream.set_position(saved_pos);
439            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
440            let expected = &[".m64n104k16", ".m64n112k16", ".m64n120k16", ".m64n128k16", ".m64n136k16", ".m64n144k16", ".m64n152k16", ".m64n160k16", ".m64n168k16", ".m64n176k16", ".m64n184k16", ".m64n192k16", ".m64n200k16", ".m64n208k16", ".m64n216k16", ".m64n224k16", ".m64n232k16", ".m64n240k16", ".m64n248k16", ".m64n256k16", ".m64n16k16", ".m64n24k16", ".m64n32k16", ".m64n40k16", ".m64n48k16", ".m64n56k16", ".m64n64k16", ".m64n72k16", ".m64n80k16", ".m64n88k16", ".m64n96k16", ".m64n8k16"];
441            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
442            Err(crate::parser::unexpected_value(span, expected, found))
443        }
444    }
445
446    impl PtxParser for WgmmaMmaAsyncSyncAlignedShapeDtypeF16F16 {
447        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
448            stream.expect_string("wgmma")?;
449            stream.expect_string(".mma_async")?;
450            let mma_async = ();
451            stream.expect_complete()?;
452            stream.expect_string(".sync")?;
453            let sync = ();
454            stream.expect_complete()?;
455            stream.expect_string(".aligned")?;
456            let aligned = ();
457            stream.expect_complete()?;
458            let shape = Shape::parse(stream)?;
459            stream.expect_complete()?;
460            let dtype = Dtype::parse(stream)?;
461            stream.expect_complete()?;
462            stream.expect_string(".f16")?;
463            let f16 = ();
464            stream.expect_complete()?;
465            stream.expect_string(".f16")?;
466            let f162 = ();
467            stream.expect_complete()?;
468            let d = GeneralOperand::parse(stream)?;
469            stream.expect_complete()?;
470            stream.expect(&PtxToken::Comma)?;
471            let a_desc = GeneralOperand::parse(stream)?;
472            stream.expect_complete()?;
473            stream.expect(&PtxToken::Comma)?;
474            let b_desc = GeneralOperand::parse(stream)?;
475            stream.expect_complete()?;
476            stream.expect(&PtxToken::Comma)?;
477            let scale_d = GeneralOperand::parse(stream)?;
478            stream.expect_complete()?;
479            stream.expect(&PtxToken::Comma)?;
480            let imm_scale_a = GeneralOperand::parse(stream)?;
481            stream.expect_complete()?;
482            stream.expect(&PtxToken::Comma)?;
483            let imm_scale_b = GeneralOperand::parse(stream)?;
484            stream.expect_complete()?;
485            stream.expect(&PtxToken::Comma)?;
486            let imm_trans_a = GeneralOperand::parse(stream)?;
487            stream.expect_complete()?;
488            stream.expect(&PtxToken::Comma)?;
489            let imm_trans_b = GeneralOperand::parse(stream)?;
490            stream.expect_complete()?;
491            stream.expect_complete()?;
492            stream.expect(&PtxToken::Semicolon)?;
493            Ok(WgmmaMmaAsyncSyncAlignedShapeDtypeF16F16 {
494                mma_async,
495                sync,
496                aligned,
497                shape,
498                dtype,
499                f16,
500                f162,
501                d,
502                a_desc,
503                b_desc,
504                scale_d,
505                imm_scale_a,
506                imm_scale_b,
507                imm_trans_a,
508                imm_trans_b,
509            })
510        }
511    }
512
513
514    impl PtxParser for WgmmaMmaAsyncSyncAlignedShapeDtypeF16F161 {
515        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
516            stream.expect_string("wgmma")?;
517            stream.expect_string(".mma_async")?;
518            let mma_async = ();
519            stream.expect_complete()?;
520            stream.expect_string(".sync")?;
521            let sync = ();
522            stream.expect_complete()?;
523            stream.expect_string(".aligned")?;
524            let aligned = ();
525            stream.expect_complete()?;
526            let shape = Shape::parse(stream)?;
527            stream.expect_complete()?;
528            let dtype = Dtype::parse(stream)?;
529            stream.expect_complete()?;
530            stream.expect_string(".f16")?;
531            let f16 = ();
532            stream.expect_complete()?;
533            stream.expect_string(".f16")?;
534            let f162 = ();
535            stream.expect_complete()?;
536            let d = GeneralOperand::parse(stream)?;
537            stream.expect_complete()?;
538            stream.expect(&PtxToken::Comma)?;
539            let a = GeneralOperand::parse(stream)?;
540            stream.expect_complete()?;
541            stream.expect(&PtxToken::Comma)?;
542            let b_desc = GeneralOperand::parse(stream)?;
543            stream.expect_complete()?;
544            stream.expect(&PtxToken::Comma)?;
545            let scale_d = GeneralOperand::parse(stream)?;
546            stream.expect_complete()?;
547            stream.expect(&PtxToken::Comma)?;
548            let imm_scale_a = GeneralOperand::parse(stream)?;
549            stream.expect_complete()?;
550            stream.expect(&PtxToken::Comma)?;
551            let imm_scale_b = GeneralOperand::parse(stream)?;
552            stream.expect_complete()?;
553            stream.expect(&PtxToken::Comma)?;
554            let imm_trans_b = GeneralOperand::parse(stream)?;
555            stream.expect_complete()?;
556            stream.expect_complete()?;
557            stream.expect(&PtxToken::Semicolon)?;
558            Ok(WgmmaMmaAsyncSyncAlignedShapeDtypeF16F161 {
559                mma_async,
560                sync,
561                aligned,
562                shape,
563                dtype,
564                f16,
565                f162,
566                d,
567                a,
568                b_desc,
569                scale_d,
570                imm_scale_a,
571                imm_scale_b,
572                imm_trans_b,
573            })
574        }
575    }
576
577
578}
579
580pub mod section_1 {
581    use super::*;
582    use crate::r#type::instruction::wgmma_mma_async::section_1::*;
583
584    // ============================================================================
585    // Generated enum parsers
586    // ============================================================================
587
588    impl PtxParser for Dtype {
589        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
590            // Try F32
591            {
592                let saved_pos = stream.position();
593                if stream.expect_string(".f32").is_ok() {
594                    return Ok(Dtype::F32);
595                }
596                stream.set_position(saved_pos);
597            }
598            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
599            let expected = &[".f32"];
600            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
601            Err(crate::parser::unexpected_value(span, expected, found))
602        }
603    }
604
605    impl PtxParser for Shape {
606        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
607            // Try M64n104k16
608            {
609                let saved_pos = stream.position();
610                if stream.expect_string(".m64n104k16").is_ok() {
611                    return Ok(Shape::M64n104k16);
612                }
613                stream.set_position(saved_pos);
614            }
615            let saved_pos = stream.position();
616            // Try M64n112k16
617            {
618                let saved_pos = stream.position();
619                if stream.expect_string(".m64n112k16").is_ok() {
620                    return Ok(Shape::M64n112k16);
621                }
622                stream.set_position(saved_pos);
623            }
624            stream.set_position(saved_pos);
625            let saved_pos = stream.position();
626            // Try M64n120k16
627            {
628                let saved_pos = stream.position();
629                if stream.expect_string(".m64n120k16").is_ok() {
630                    return Ok(Shape::M64n120k16);
631                }
632                stream.set_position(saved_pos);
633            }
634            stream.set_position(saved_pos);
635            let saved_pos = stream.position();
636            // Try M64n128k16
637            {
638                let saved_pos = stream.position();
639                if stream.expect_string(".m64n128k16").is_ok() {
640                    return Ok(Shape::M64n128k16);
641                }
642                stream.set_position(saved_pos);
643            }
644            stream.set_position(saved_pos);
645            let saved_pos = stream.position();
646            // Try M64n136k16
647            {
648                let saved_pos = stream.position();
649                if stream.expect_string(".m64n136k16").is_ok() {
650                    return Ok(Shape::M64n136k16);
651                }
652                stream.set_position(saved_pos);
653            }
654            stream.set_position(saved_pos);
655            let saved_pos = stream.position();
656            // Try M64n144k16
657            {
658                let saved_pos = stream.position();
659                if stream.expect_string(".m64n144k16").is_ok() {
660                    return Ok(Shape::M64n144k16);
661                }
662                stream.set_position(saved_pos);
663            }
664            stream.set_position(saved_pos);
665            let saved_pos = stream.position();
666            // Try M64n152k16
667            {
668                let saved_pos = stream.position();
669                if stream.expect_string(".m64n152k16").is_ok() {
670                    return Ok(Shape::M64n152k16);
671                }
672                stream.set_position(saved_pos);
673            }
674            stream.set_position(saved_pos);
675            let saved_pos = stream.position();
676            // Try M64n160k16
677            {
678                let saved_pos = stream.position();
679                if stream.expect_string(".m64n160k16").is_ok() {
680                    return Ok(Shape::M64n160k16);
681                }
682                stream.set_position(saved_pos);
683            }
684            stream.set_position(saved_pos);
685            let saved_pos = stream.position();
686            // Try M64n168k16
687            {
688                let saved_pos = stream.position();
689                if stream.expect_string(".m64n168k16").is_ok() {
690                    return Ok(Shape::M64n168k16);
691                }
692                stream.set_position(saved_pos);
693            }
694            stream.set_position(saved_pos);
695            let saved_pos = stream.position();
696            // Try M64n176k16
697            {
698                let saved_pos = stream.position();
699                if stream.expect_string(".m64n176k16").is_ok() {
700                    return Ok(Shape::M64n176k16);
701                }
702                stream.set_position(saved_pos);
703            }
704            stream.set_position(saved_pos);
705            let saved_pos = stream.position();
706            // Try M64n184k16
707            {
708                let saved_pos = stream.position();
709                if stream.expect_string(".m64n184k16").is_ok() {
710                    return Ok(Shape::M64n184k16);
711                }
712                stream.set_position(saved_pos);
713            }
714            stream.set_position(saved_pos);
715            let saved_pos = stream.position();
716            // Try M64n192k16
717            {
718                let saved_pos = stream.position();
719                if stream.expect_string(".m64n192k16").is_ok() {
720                    return Ok(Shape::M64n192k16);
721                }
722                stream.set_position(saved_pos);
723            }
724            stream.set_position(saved_pos);
725            let saved_pos = stream.position();
726            // Try M64n200k16
727            {
728                let saved_pos = stream.position();
729                if stream.expect_string(".m64n200k16").is_ok() {
730                    return Ok(Shape::M64n200k16);
731                }
732                stream.set_position(saved_pos);
733            }
734            stream.set_position(saved_pos);
735            let saved_pos = stream.position();
736            // Try M64n208k16
737            {
738                let saved_pos = stream.position();
739                if stream.expect_string(".m64n208k16").is_ok() {
740                    return Ok(Shape::M64n208k16);
741                }
742                stream.set_position(saved_pos);
743            }
744            stream.set_position(saved_pos);
745            let saved_pos = stream.position();
746            // Try M64n216k16
747            {
748                let saved_pos = stream.position();
749                if stream.expect_string(".m64n216k16").is_ok() {
750                    return Ok(Shape::M64n216k16);
751                }
752                stream.set_position(saved_pos);
753            }
754            stream.set_position(saved_pos);
755            let saved_pos = stream.position();
756            // Try M64n224k16
757            {
758                let saved_pos = stream.position();
759                if stream.expect_string(".m64n224k16").is_ok() {
760                    return Ok(Shape::M64n224k16);
761                }
762                stream.set_position(saved_pos);
763            }
764            stream.set_position(saved_pos);
765            let saved_pos = stream.position();
766            // Try M64n232k16
767            {
768                let saved_pos = stream.position();
769                if stream.expect_string(".m64n232k16").is_ok() {
770                    return Ok(Shape::M64n232k16);
771                }
772                stream.set_position(saved_pos);
773            }
774            stream.set_position(saved_pos);
775            let saved_pos = stream.position();
776            // Try M64n240k16
777            {
778                let saved_pos = stream.position();
779                if stream.expect_string(".m64n240k16").is_ok() {
780                    return Ok(Shape::M64n240k16);
781                }
782                stream.set_position(saved_pos);
783            }
784            stream.set_position(saved_pos);
785            let saved_pos = stream.position();
786            // Try M64n248k16
787            {
788                let saved_pos = stream.position();
789                if stream.expect_string(".m64n248k16").is_ok() {
790                    return Ok(Shape::M64n248k16);
791                }
792                stream.set_position(saved_pos);
793            }
794            stream.set_position(saved_pos);
795            let saved_pos = stream.position();
796            // Try M64n256k16
797            {
798                let saved_pos = stream.position();
799                if stream.expect_string(".m64n256k16").is_ok() {
800                    return Ok(Shape::M64n256k16);
801                }
802                stream.set_position(saved_pos);
803            }
804            stream.set_position(saved_pos);
805            let saved_pos = stream.position();
806            // Try M64n16k16
807            {
808                let saved_pos = stream.position();
809                if stream.expect_string(".m64n16k16").is_ok() {
810                    return Ok(Shape::M64n16k16);
811                }
812                stream.set_position(saved_pos);
813            }
814            stream.set_position(saved_pos);
815            let saved_pos = stream.position();
816            // Try M64n24k16
817            {
818                let saved_pos = stream.position();
819                if stream.expect_string(".m64n24k16").is_ok() {
820                    return Ok(Shape::M64n24k16);
821                }
822                stream.set_position(saved_pos);
823            }
824            stream.set_position(saved_pos);
825            let saved_pos = stream.position();
826            // Try M64n32k16
827            {
828                let saved_pos = stream.position();
829                if stream.expect_string(".m64n32k16").is_ok() {
830                    return Ok(Shape::M64n32k16);
831                }
832                stream.set_position(saved_pos);
833            }
834            stream.set_position(saved_pos);
835            let saved_pos = stream.position();
836            // Try M64n40k16
837            {
838                let saved_pos = stream.position();
839                if stream.expect_string(".m64n40k16").is_ok() {
840                    return Ok(Shape::M64n40k16);
841                }
842                stream.set_position(saved_pos);
843            }
844            stream.set_position(saved_pos);
845            let saved_pos = stream.position();
846            // Try M64n48k16
847            {
848                let saved_pos = stream.position();
849                if stream.expect_string(".m64n48k16").is_ok() {
850                    return Ok(Shape::M64n48k16);
851                }
852                stream.set_position(saved_pos);
853            }
854            stream.set_position(saved_pos);
855            let saved_pos = stream.position();
856            // Try M64n56k16
857            {
858                let saved_pos = stream.position();
859                if stream.expect_string(".m64n56k16").is_ok() {
860                    return Ok(Shape::M64n56k16);
861                }
862                stream.set_position(saved_pos);
863            }
864            stream.set_position(saved_pos);
865            let saved_pos = stream.position();
866            // Try M64n64k16
867            {
868                let saved_pos = stream.position();
869                if stream.expect_string(".m64n64k16").is_ok() {
870                    return Ok(Shape::M64n64k16);
871                }
872                stream.set_position(saved_pos);
873            }
874            stream.set_position(saved_pos);
875            let saved_pos = stream.position();
876            // Try M64n72k16
877            {
878                let saved_pos = stream.position();
879                if stream.expect_string(".m64n72k16").is_ok() {
880                    return Ok(Shape::M64n72k16);
881                }
882                stream.set_position(saved_pos);
883            }
884            stream.set_position(saved_pos);
885            let saved_pos = stream.position();
886            // Try M64n80k16
887            {
888                let saved_pos = stream.position();
889                if stream.expect_string(".m64n80k16").is_ok() {
890                    return Ok(Shape::M64n80k16);
891                }
892                stream.set_position(saved_pos);
893            }
894            stream.set_position(saved_pos);
895            let saved_pos = stream.position();
896            // Try M64n88k16
897            {
898                let saved_pos = stream.position();
899                if stream.expect_string(".m64n88k16").is_ok() {
900                    return Ok(Shape::M64n88k16);
901                }
902                stream.set_position(saved_pos);
903            }
904            stream.set_position(saved_pos);
905            let saved_pos = stream.position();
906            // Try M64n96k16
907            {
908                let saved_pos = stream.position();
909                if stream.expect_string(".m64n96k16").is_ok() {
910                    return Ok(Shape::M64n96k16);
911                }
912                stream.set_position(saved_pos);
913            }
914            stream.set_position(saved_pos);
915            let saved_pos = stream.position();
916            // Try M64n8k16
917            {
918                let saved_pos = stream.position();
919                if stream.expect_string(".m64n8k16").is_ok() {
920                    return Ok(Shape::M64n8k16);
921                }
922                stream.set_position(saved_pos);
923            }
924            stream.set_position(saved_pos);
925            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
926            let expected = &[".m64n104k16", ".m64n112k16", ".m64n120k16", ".m64n128k16", ".m64n136k16", ".m64n144k16", ".m64n152k16", ".m64n160k16", ".m64n168k16", ".m64n176k16", ".m64n184k16", ".m64n192k16", ".m64n200k16", ".m64n208k16", ".m64n216k16", ".m64n224k16", ".m64n232k16", ".m64n240k16", ".m64n248k16", ".m64n256k16", ".m64n16k16", ".m64n24k16", ".m64n32k16", ".m64n40k16", ".m64n48k16", ".m64n56k16", ".m64n64k16", ".m64n72k16", ".m64n80k16", ".m64n88k16", ".m64n96k16", ".m64n8k16"];
927            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
928            Err(crate::parser::unexpected_value(span, expected, found))
929        }
930    }
931
932    impl PtxParser for WgmmaMmaAsyncSyncAlignedShapeDtypeBf16Bf16 {
933        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
934            stream.expect_string("wgmma")?;
935            stream.expect_string(".mma_async")?;
936            let mma_async = ();
937            stream.expect_complete()?;
938            stream.expect_string(".sync")?;
939            let sync = ();
940            stream.expect_complete()?;
941            stream.expect_string(".aligned")?;
942            let aligned = ();
943            stream.expect_complete()?;
944            let shape = Shape::parse(stream)?;
945            stream.expect_complete()?;
946            let dtype = Dtype::parse(stream)?;
947            stream.expect_complete()?;
948            stream.expect_string(".bf16")?;
949            let bf16 = ();
950            stream.expect_complete()?;
951            stream.expect_string(".bf16")?;
952            let bf162 = ();
953            stream.expect_complete()?;
954            let d = GeneralOperand::parse(stream)?;
955            stream.expect_complete()?;
956            stream.expect(&PtxToken::Comma)?;
957            let a_desc = GeneralOperand::parse(stream)?;
958            stream.expect_complete()?;
959            stream.expect(&PtxToken::Comma)?;
960            let b_desc = GeneralOperand::parse(stream)?;
961            stream.expect_complete()?;
962            stream.expect(&PtxToken::Comma)?;
963            let scale_d = GeneralOperand::parse(stream)?;
964            stream.expect_complete()?;
965            stream.expect(&PtxToken::Comma)?;
966            let imm_scale_a = GeneralOperand::parse(stream)?;
967            stream.expect_complete()?;
968            stream.expect(&PtxToken::Comma)?;
969            let imm_scale_b = GeneralOperand::parse(stream)?;
970            stream.expect_complete()?;
971            stream.expect(&PtxToken::Comma)?;
972            let imm_trans_a = GeneralOperand::parse(stream)?;
973            stream.expect_complete()?;
974            stream.expect(&PtxToken::Comma)?;
975            let imm_trans_b = GeneralOperand::parse(stream)?;
976            stream.expect_complete()?;
977            stream.expect_complete()?;
978            stream.expect(&PtxToken::Semicolon)?;
979            Ok(WgmmaMmaAsyncSyncAlignedShapeDtypeBf16Bf16 {
980                mma_async,
981                sync,
982                aligned,
983                shape,
984                dtype,
985                bf16,
986                bf162,
987                d,
988                a_desc,
989                b_desc,
990                scale_d,
991                imm_scale_a,
992                imm_scale_b,
993                imm_trans_a,
994                imm_trans_b,
995            })
996        }
997    }
998
999
1000    impl PtxParser for WgmmaMmaAsyncSyncAlignedShapeDtypeBf16Bf161 {
1001        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1002            stream.expect_string("wgmma")?;
1003            stream.expect_string(".mma_async")?;
1004            let mma_async = ();
1005            stream.expect_complete()?;
1006            stream.expect_string(".sync")?;
1007            let sync = ();
1008            stream.expect_complete()?;
1009            stream.expect_string(".aligned")?;
1010            let aligned = ();
1011            stream.expect_complete()?;
1012            let shape = Shape::parse(stream)?;
1013            stream.expect_complete()?;
1014            let dtype = Dtype::parse(stream)?;
1015            stream.expect_complete()?;
1016            stream.expect_string(".bf16")?;
1017            let bf16 = ();
1018            stream.expect_complete()?;
1019            stream.expect_string(".bf16")?;
1020            let bf162 = ();
1021            stream.expect_complete()?;
1022            let d = GeneralOperand::parse(stream)?;
1023            stream.expect_complete()?;
1024            stream.expect(&PtxToken::Comma)?;
1025            let a = GeneralOperand::parse(stream)?;
1026            stream.expect_complete()?;
1027            stream.expect(&PtxToken::Comma)?;
1028            let b_desc = GeneralOperand::parse(stream)?;
1029            stream.expect_complete()?;
1030            stream.expect(&PtxToken::Comma)?;
1031            let scale_d = GeneralOperand::parse(stream)?;
1032            stream.expect_complete()?;
1033            stream.expect(&PtxToken::Comma)?;
1034            let imm_scale_a = GeneralOperand::parse(stream)?;
1035            stream.expect_complete()?;
1036            stream.expect(&PtxToken::Comma)?;
1037            let imm_scale_b = GeneralOperand::parse(stream)?;
1038            stream.expect_complete()?;
1039            stream.expect(&PtxToken::Comma)?;
1040            let imm_trans_b = GeneralOperand::parse(stream)?;
1041            stream.expect_complete()?;
1042            stream.expect_complete()?;
1043            stream.expect(&PtxToken::Semicolon)?;
1044            Ok(WgmmaMmaAsyncSyncAlignedShapeDtypeBf16Bf161 {
1045                mma_async,
1046                sync,
1047                aligned,
1048                shape,
1049                dtype,
1050                bf16,
1051                bf162,
1052                d,
1053                a,
1054                b_desc,
1055                scale_d,
1056                imm_scale_a,
1057                imm_scale_b,
1058                imm_trans_b,
1059            })
1060        }
1061    }
1062
1063
1064}
1065
1066pub mod section_2 {
1067    use super::*;
1068    use crate::r#type::instruction::wgmma_mma_async::section_2::*;
1069
1070    // ============================================================================
1071    // Generated enum parsers
1072    // ============================================================================
1073
1074    impl PtxParser for Dtype {
1075        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1076            // Try F32
1077            {
1078                let saved_pos = stream.position();
1079                if stream.expect_string(".f32").is_ok() {
1080                    return Ok(Dtype::F32);
1081                }
1082                stream.set_position(saved_pos);
1083            }
1084            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1085            let expected = &[".f32"];
1086            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1087            Err(crate::parser::unexpected_value(span, expected, found))
1088        }
1089    }
1090
1091    impl PtxParser for Shape {
1092        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1093            // Try M64n104k8
1094            {
1095                let saved_pos = stream.position();
1096                if stream.expect_string(".m64n104k8").is_ok() {
1097                    return Ok(Shape::M64n104k8);
1098                }
1099                stream.set_position(saved_pos);
1100            }
1101            let saved_pos = stream.position();
1102            // Try M64n112k8
1103            {
1104                let saved_pos = stream.position();
1105                if stream.expect_string(".m64n112k8").is_ok() {
1106                    return Ok(Shape::M64n112k8);
1107                }
1108                stream.set_position(saved_pos);
1109            }
1110            stream.set_position(saved_pos);
1111            let saved_pos = stream.position();
1112            // Try M64n120k8
1113            {
1114                let saved_pos = stream.position();
1115                if stream.expect_string(".m64n120k8").is_ok() {
1116                    return Ok(Shape::M64n120k8);
1117                }
1118                stream.set_position(saved_pos);
1119            }
1120            stream.set_position(saved_pos);
1121            let saved_pos = stream.position();
1122            // Try M64n128k8
1123            {
1124                let saved_pos = stream.position();
1125                if stream.expect_string(".m64n128k8").is_ok() {
1126                    return Ok(Shape::M64n128k8);
1127                }
1128                stream.set_position(saved_pos);
1129            }
1130            stream.set_position(saved_pos);
1131            let saved_pos = stream.position();
1132            // Try M64n136k8
1133            {
1134                let saved_pos = stream.position();
1135                if stream.expect_string(".m64n136k8").is_ok() {
1136                    return Ok(Shape::M64n136k8);
1137                }
1138                stream.set_position(saved_pos);
1139            }
1140            stream.set_position(saved_pos);
1141            let saved_pos = stream.position();
1142            // Try M64n144k8
1143            {
1144                let saved_pos = stream.position();
1145                if stream.expect_string(".m64n144k8").is_ok() {
1146                    return Ok(Shape::M64n144k8);
1147                }
1148                stream.set_position(saved_pos);
1149            }
1150            stream.set_position(saved_pos);
1151            let saved_pos = stream.position();
1152            // Try M64n152k8
1153            {
1154                let saved_pos = stream.position();
1155                if stream.expect_string(".m64n152k8").is_ok() {
1156                    return Ok(Shape::M64n152k8);
1157                }
1158                stream.set_position(saved_pos);
1159            }
1160            stream.set_position(saved_pos);
1161            let saved_pos = stream.position();
1162            // Try M64n160k8
1163            {
1164                let saved_pos = stream.position();
1165                if stream.expect_string(".m64n160k8").is_ok() {
1166                    return Ok(Shape::M64n160k8);
1167                }
1168                stream.set_position(saved_pos);
1169            }
1170            stream.set_position(saved_pos);
1171            let saved_pos = stream.position();
1172            // Try M64n168k8
1173            {
1174                let saved_pos = stream.position();
1175                if stream.expect_string(".m64n168k8").is_ok() {
1176                    return Ok(Shape::M64n168k8);
1177                }
1178                stream.set_position(saved_pos);
1179            }
1180            stream.set_position(saved_pos);
1181            let saved_pos = stream.position();
1182            // Try M64n176k8
1183            {
1184                let saved_pos = stream.position();
1185                if stream.expect_string(".m64n176k8").is_ok() {
1186                    return Ok(Shape::M64n176k8);
1187                }
1188                stream.set_position(saved_pos);
1189            }
1190            stream.set_position(saved_pos);
1191            let saved_pos = stream.position();
1192            // Try M64n184k8
1193            {
1194                let saved_pos = stream.position();
1195                if stream.expect_string(".m64n184k8").is_ok() {
1196                    return Ok(Shape::M64n184k8);
1197                }
1198                stream.set_position(saved_pos);
1199            }
1200            stream.set_position(saved_pos);
1201            let saved_pos = stream.position();
1202            // Try M64n192k8
1203            {
1204                let saved_pos = stream.position();
1205                if stream.expect_string(".m64n192k8").is_ok() {
1206                    return Ok(Shape::M64n192k8);
1207                }
1208                stream.set_position(saved_pos);
1209            }
1210            stream.set_position(saved_pos);
1211            let saved_pos = stream.position();
1212            // Try M64n200k8
1213            {
1214                let saved_pos = stream.position();
1215                if stream.expect_string(".m64n200k8").is_ok() {
1216                    return Ok(Shape::M64n200k8);
1217                }
1218                stream.set_position(saved_pos);
1219            }
1220            stream.set_position(saved_pos);
1221            let saved_pos = stream.position();
1222            // Try M64n208k8
1223            {
1224                let saved_pos = stream.position();
1225                if stream.expect_string(".m64n208k8").is_ok() {
1226                    return Ok(Shape::M64n208k8);
1227                }
1228                stream.set_position(saved_pos);
1229            }
1230            stream.set_position(saved_pos);
1231            let saved_pos = stream.position();
1232            // Try M64n216k8
1233            {
1234                let saved_pos = stream.position();
1235                if stream.expect_string(".m64n216k8").is_ok() {
1236                    return Ok(Shape::M64n216k8);
1237                }
1238                stream.set_position(saved_pos);
1239            }
1240            stream.set_position(saved_pos);
1241            let saved_pos = stream.position();
1242            // Try M64n224k8
1243            {
1244                let saved_pos = stream.position();
1245                if stream.expect_string(".m64n224k8").is_ok() {
1246                    return Ok(Shape::M64n224k8);
1247                }
1248                stream.set_position(saved_pos);
1249            }
1250            stream.set_position(saved_pos);
1251            let saved_pos = stream.position();
1252            // Try M64n232k8
1253            {
1254                let saved_pos = stream.position();
1255                if stream.expect_string(".m64n232k8").is_ok() {
1256                    return Ok(Shape::M64n232k8);
1257                }
1258                stream.set_position(saved_pos);
1259            }
1260            stream.set_position(saved_pos);
1261            let saved_pos = stream.position();
1262            // Try M64n240k8
1263            {
1264                let saved_pos = stream.position();
1265                if stream.expect_string(".m64n240k8").is_ok() {
1266                    return Ok(Shape::M64n240k8);
1267                }
1268                stream.set_position(saved_pos);
1269            }
1270            stream.set_position(saved_pos);
1271            let saved_pos = stream.position();
1272            // Try M64n248k8
1273            {
1274                let saved_pos = stream.position();
1275                if stream.expect_string(".m64n248k8").is_ok() {
1276                    return Ok(Shape::M64n248k8);
1277                }
1278                stream.set_position(saved_pos);
1279            }
1280            stream.set_position(saved_pos);
1281            let saved_pos = stream.position();
1282            // Try M64n256k8
1283            {
1284                let saved_pos = stream.position();
1285                if stream.expect_string(".m64n256k8").is_ok() {
1286                    return Ok(Shape::M64n256k8);
1287                }
1288                stream.set_position(saved_pos);
1289            }
1290            stream.set_position(saved_pos);
1291            let saved_pos = stream.position();
1292            // Try M64n16k8
1293            {
1294                let saved_pos = stream.position();
1295                if stream.expect_string(".m64n16k8").is_ok() {
1296                    return Ok(Shape::M64n16k8);
1297                }
1298                stream.set_position(saved_pos);
1299            }
1300            stream.set_position(saved_pos);
1301            let saved_pos = stream.position();
1302            // Try M64n24k8
1303            {
1304                let saved_pos = stream.position();
1305                if stream.expect_string(".m64n24k8").is_ok() {
1306                    return Ok(Shape::M64n24k8);
1307                }
1308                stream.set_position(saved_pos);
1309            }
1310            stream.set_position(saved_pos);
1311            let saved_pos = stream.position();
1312            // Try M64n32k8
1313            {
1314                let saved_pos = stream.position();
1315                if stream.expect_string(".m64n32k8").is_ok() {
1316                    return Ok(Shape::M64n32k8);
1317                }
1318                stream.set_position(saved_pos);
1319            }
1320            stream.set_position(saved_pos);
1321            let saved_pos = stream.position();
1322            // Try M64n40k8
1323            {
1324                let saved_pos = stream.position();
1325                if stream.expect_string(".m64n40k8").is_ok() {
1326                    return Ok(Shape::M64n40k8);
1327                }
1328                stream.set_position(saved_pos);
1329            }
1330            stream.set_position(saved_pos);
1331            let saved_pos = stream.position();
1332            // Try M64n48k8
1333            {
1334                let saved_pos = stream.position();
1335                if stream.expect_string(".m64n48k8").is_ok() {
1336                    return Ok(Shape::M64n48k8);
1337                }
1338                stream.set_position(saved_pos);
1339            }
1340            stream.set_position(saved_pos);
1341            let saved_pos = stream.position();
1342            // Try M64n56k8
1343            {
1344                let saved_pos = stream.position();
1345                if stream.expect_string(".m64n56k8").is_ok() {
1346                    return Ok(Shape::M64n56k8);
1347                }
1348                stream.set_position(saved_pos);
1349            }
1350            stream.set_position(saved_pos);
1351            let saved_pos = stream.position();
1352            // Try M64n64k8
1353            {
1354                let saved_pos = stream.position();
1355                if stream.expect_string(".m64n64k8").is_ok() {
1356                    return Ok(Shape::M64n64k8);
1357                }
1358                stream.set_position(saved_pos);
1359            }
1360            stream.set_position(saved_pos);
1361            let saved_pos = stream.position();
1362            // Try M64n72k8
1363            {
1364                let saved_pos = stream.position();
1365                if stream.expect_string(".m64n72k8").is_ok() {
1366                    return Ok(Shape::M64n72k8);
1367                }
1368                stream.set_position(saved_pos);
1369            }
1370            stream.set_position(saved_pos);
1371            let saved_pos = stream.position();
1372            // Try M64n80k8
1373            {
1374                let saved_pos = stream.position();
1375                if stream.expect_string(".m64n80k8").is_ok() {
1376                    return Ok(Shape::M64n80k8);
1377                }
1378                stream.set_position(saved_pos);
1379            }
1380            stream.set_position(saved_pos);
1381            let saved_pos = stream.position();
1382            // Try M64n88k8
1383            {
1384                let saved_pos = stream.position();
1385                if stream.expect_string(".m64n88k8").is_ok() {
1386                    return Ok(Shape::M64n88k8);
1387                }
1388                stream.set_position(saved_pos);
1389            }
1390            stream.set_position(saved_pos);
1391            let saved_pos = stream.position();
1392            // Try M64n96k8
1393            {
1394                let saved_pos = stream.position();
1395                if stream.expect_string(".m64n96k8").is_ok() {
1396                    return Ok(Shape::M64n96k8);
1397                }
1398                stream.set_position(saved_pos);
1399            }
1400            stream.set_position(saved_pos);
1401            let saved_pos = stream.position();
1402            // Try M64n8k8
1403            {
1404                let saved_pos = stream.position();
1405                if stream.expect_string(".m64n8k8").is_ok() {
1406                    return Ok(Shape::M64n8k8);
1407                }
1408                stream.set_position(saved_pos);
1409            }
1410            stream.set_position(saved_pos);
1411            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1412            let expected = &[".m64n104k8", ".m64n112k8", ".m64n120k8", ".m64n128k8", ".m64n136k8", ".m64n144k8", ".m64n152k8", ".m64n160k8", ".m64n168k8", ".m64n176k8", ".m64n184k8", ".m64n192k8", ".m64n200k8", ".m64n208k8", ".m64n216k8", ".m64n224k8", ".m64n232k8", ".m64n240k8", ".m64n248k8", ".m64n256k8", ".m64n16k8", ".m64n24k8", ".m64n32k8", ".m64n40k8", ".m64n48k8", ".m64n56k8", ".m64n64k8", ".m64n72k8", ".m64n80k8", ".m64n88k8", ".m64n96k8", ".m64n8k8"];
1413            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1414            Err(crate::parser::unexpected_value(span, expected, found))
1415        }
1416    }
1417
1418    impl PtxParser for WgmmaMmaAsyncSyncAlignedShapeDtypeTf32Tf32 {
1419        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1420            stream.expect_string("wgmma")?;
1421            stream.expect_string(".mma_async")?;
1422            let mma_async = ();
1423            stream.expect_complete()?;
1424            stream.expect_string(".sync")?;
1425            let sync = ();
1426            stream.expect_complete()?;
1427            stream.expect_string(".aligned")?;
1428            let aligned = ();
1429            stream.expect_complete()?;
1430            let shape = Shape::parse(stream)?;
1431            stream.expect_complete()?;
1432            let dtype = Dtype::parse(stream)?;
1433            stream.expect_complete()?;
1434            stream.expect_string(".tf32")?;
1435            let tf32 = ();
1436            stream.expect_complete()?;
1437            stream.expect_string(".tf32")?;
1438            let tf322 = ();
1439            stream.expect_complete()?;
1440            let d = GeneralOperand::parse(stream)?;
1441            stream.expect_complete()?;
1442            stream.expect(&PtxToken::Comma)?;
1443            let a_desc = GeneralOperand::parse(stream)?;
1444            stream.expect_complete()?;
1445            stream.expect(&PtxToken::Comma)?;
1446            let b_desc = GeneralOperand::parse(stream)?;
1447            stream.expect_complete()?;
1448            stream.expect(&PtxToken::Comma)?;
1449            let scale_d = GeneralOperand::parse(stream)?;
1450            stream.expect_complete()?;
1451            stream.expect(&PtxToken::Comma)?;
1452            let imm_scale_a = GeneralOperand::parse(stream)?;
1453            stream.expect_complete()?;
1454            stream.expect(&PtxToken::Comma)?;
1455            let imm_scale_b = GeneralOperand::parse(stream)?;
1456            stream.expect_complete()?;
1457            stream.expect_complete()?;
1458            stream.expect(&PtxToken::Semicolon)?;
1459            Ok(WgmmaMmaAsyncSyncAlignedShapeDtypeTf32Tf32 {
1460                mma_async,
1461                sync,
1462                aligned,
1463                shape,
1464                dtype,
1465                tf32,
1466                tf322,
1467                d,
1468                a_desc,
1469                b_desc,
1470                scale_d,
1471                imm_scale_a,
1472                imm_scale_b,
1473            })
1474        }
1475    }
1476
1477
1478    impl PtxParser for WgmmaMmaAsyncSyncAlignedShapeDtypeTf32Tf321 {
1479        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1480            stream.expect_string("wgmma")?;
1481            stream.expect_string(".mma_async")?;
1482            let mma_async = ();
1483            stream.expect_complete()?;
1484            stream.expect_string(".sync")?;
1485            let sync = ();
1486            stream.expect_complete()?;
1487            stream.expect_string(".aligned")?;
1488            let aligned = ();
1489            stream.expect_complete()?;
1490            let shape = Shape::parse(stream)?;
1491            stream.expect_complete()?;
1492            let dtype = Dtype::parse(stream)?;
1493            stream.expect_complete()?;
1494            stream.expect_string(".tf32")?;
1495            let tf32 = ();
1496            stream.expect_complete()?;
1497            stream.expect_string(".tf32")?;
1498            let tf322 = ();
1499            stream.expect_complete()?;
1500            let d = GeneralOperand::parse(stream)?;
1501            stream.expect_complete()?;
1502            stream.expect(&PtxToken::Comma)?;
1503            let a = GeneralOperand::parse(stream)?;
1504            stream.expect_complete()?;
1505            stream.expect(&PtxToken::Comma)?;
1506            let b_desc = GeneralOperand::parse(stream)?;
1507            stream.expect_complete()?;
1508            stream.expect(&PtxToken::Comma)?;
1509            let scale_d = GeneralOperand::parse(stream)?;
1510            stream.expect_complete()?;
1511            stream.expect(&PtxToken::Comma)?;
1512            let imm_scale_a = GeneralOperand::parse(stream)?;
1513            stream.expect_complete()?;
1514            stream.expect(&PtxToken::Comma)?;
1515            let imm_scale_b = GeneralOperand::parse(stream)?;
1516            stream.expect_complete()?;
1517            stream.expect_complete()?;
1518            stream.expect(&PtxToken::Semicolon)?;
1519            Ok(WgmmaMmaAsyncSyncAlignedShapeDtypeTf32Tf321 {
1520                mma_async,
1521                sync,
1522                aligned,
1523                shape,
1524                dtype,
1525                tf32,
1526                tf322,
1527                d,
1528                a,
1529                b_desc,
1530                scale_d,
1531                imm_scale_a,
1532                imm_scale_b,
1533            })
1534        }
1535    }
1536
1537
1538}
1539
1540pub mod section_3 {
1541    use super::*;
1542    use crate::r#type::instruction::wgmma_mma_async::section_3::*;
1543
1544    // ============================================================================
1545    // Generated enum parsers
1546    // ============================================================================
1547
1548    impl PtxParser for Atype {
1549        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1550            // Try E4m3
1551            {
1552                let saved_pos = stream.position();
1553                if stream.expect_string(".e4m3").is_ok() {
1554                    return Ok(Atype::E4m3);
1555                }
1556                stream.set_position(saved_pos);
1557            }
1558            let saved_pos = stream.position();
1559            // Try E5m2
1560            {
1561                let saved_pos = stream.position();
1562                if stream.expect_string(".e5m2").is_ok() {
1563                    return Ok(Atype::E5m2);
1564                }
1565                stream.set_position(saved_pos);
1566            }
1567            stream.set_position(saved_pos);
1568            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1569            let expected = &[".e4m3", ".e5m2"];
1570            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1571            Err(crate::parser::unexpected_value(span, expected, found))
1572        }
1573    }
1574
1575    impl PtxParser for Btype {
1576        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1577            // Try E4m3
1578            {
1579                let saved_pos = stream.position();
1580                if stream.expect_string(".e4m3").is_ok() {
1581                    return Ok(Btype::E4m3);
1582                }
1583                stream.set_position(saved_pos);
1584            }
1585            let saved_pos = stream.position();
1586            // Try E5m2
1587            {
1588                let saved_pos = stream.position();
1589                if stream.expect_string(".e5m2").is_ok() {
1590                    return Ok(Btype::E5m2);
1591                }
1592                stream.set_position(saved_pos);
1593            }
1594            stream.set_position(saved_pos);
1595            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1596            let expected = &[".e4m3", ".e5m2"];
1597            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1598            Err(crate::parser::unexpected_value(span, expected, found))
1599        }
1600    }
1601
1602    impl PtxParser for Dtype {
1603        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1604            // Try F16
1605            {
1606                let saved_pos = stream.position();
1607                if stream.expect_string(".f16").is_ok() {
1608                    return Ok(Dtype::F16);
1609                }
1610                stream.set_position(saved_pos);
1611            }
1612            let saved_pos = stream.position();
1613            // Try F32
1614            {
1615                let saved_pos = stream.position();
1616                if stream.expect_string(".f32").is_ok() {
1617                    return Ok(Dtype::F32);
1618                }
1619                stream.set_position(saved_pos);
1620            }
1621            stream.set_position(saved_pos);
1622            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1623            let expected = &[".f16", ".f32"];
1624            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1625            Err(crate::parser::unexpected_value(span, expected, found))
1626        }
1627    }
1628
1629    impl PtxParser for Shape {
1630        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1631            // Try M64n104k32
1632            {
1633                let saved_pos = stream.position();
1634                if stream.expect_string(".m64n104k32").is_ok() {
1635                    return Ok(Shape::M64n104k32);
1636                }
1637                stream.set_position(saved_pos);
1638            }
1639            let saved_pos = stream.position();
1640            // Try M64n112k32
1641            {
1642                let saved_pos = stream.position();
1643                if stream.expect_string(".m64n112k32").is_ok() {
1644                    return Ok(Shape::M64n112k32);
1645                }
1646                stream.set_position(saved_pos);
1647            }
1648            stream.set_position(saved_pos);
1649            let saved_pos = stream.position();
1650            // Try M64n120k32
1651            {
1652                let saved_pos = stream.position();
1653                if stream.expect_string(".m64n120k32").is_ok() {
1654                    return Ok(Shape::M64n120k32);
1655                }
1656                stream.set_position(saved_pos);
1657            }
1658            stream.set_position(saved_pos);
1659            let saved_pos = stream.position();
1660            // Try M64n128k32
1661            {
1662                let saved_pos = stream.position();
1663                if stream.expect_string(".m64n128k32").is_ok() {
1664                    return Ok(Shape::M64n128k32);
1665                }
1666                stream.set_position(saved_pos);
1667            }
1668            stream.set_position(saved_pos);
1669            let saved_pos = stream.position();
1670            // Try M64n136k32
1671            {
1672                let saved_pos = stream.position();
1673                if stream.expect_string(".m64n136k32").is_ok() {
1674                    return Ok(Shape::M64n136k32);
1675                }
1676                stream.set_position(saved_pos);
1677            }
1678            stream.set_position(saved_pos);
1679            let saved_pos = stream.position();
1680            // Try M64n144k32
1681            {
1682                let saved_pos = stream.position();
1683                if stream.expect_string(".m64n144k32").is_ok() {
1684                    return Ok(Shape::M64n144k32);
1685                }
1686                stream.set_position(saved_pos);
1687            }
1688            stream.set_position(saved_pos);
1689            let saved_pos = stream.position();
1690            // Try M64n152k32
1691            {
1692                let saved_pos = stream.position();
1693                if stream.expect_string(".m64n152k32").is_ok() {
1694                    return Ok(Shape::M64n152k32);
1695                }
1696                stream.set_position(saved_pos);
1697            }
1698            stream.set_position(saved_pos);
1699            let saved_pos = stream.position();
1700            // Try M64n160k32
1701            {
1702                let saved_pos = stream.position();
1703                if stream.expect_string(".m64n160k32").is_ok() {
1704                    return Ok(Shape::M64n160k32);
1705                }
1706                stream.set_position(saved_pos);
1707            }
1708            stream.set_position(saved_pos);
1709            let saved_pos = stream.position();
1710            // Try M64n168k32
1711            {
1712                let saved_pos = stream.position();
1713                if stream.expect_string(".m64n168k32").is_ok() {
1714                    return Ok(Shape::M64n168k32);
1715                }
1716                stream.set_position(saved_pos);
1717            }
1718            stream.set_position(saved_pos);
1719            let saved_pos = stream.position();
1720            // Try M64n176k32
1721            {
1722                let saved_pos = stream.position();
1723                if stream.expect_string(".m64n176k32").is_ok() {
1724                    return Ok(Shape::M64n176k32);
1725                }
1726                stream.set_position(saved_pos);
1727            }
1728            stream.set_position(saved_pos);
1729            let saved_pos = stream.position();
1730            // Try M64n184k32
1731            {
1732                let saved_pos = stream.position();
1733                if stream.expect_string(".m64n184k32").is_ok() {
1734                    return Ok(Shape::M64n184k32);
1735                }
1736                stream.set_position(saved_pos);
1737            }
1738            stream.set_position(saved_pos);
1739            let saved_pos = stream.position();
1740            // Try M64n192k32
1741            {
1742                let saved_pos = stream.position();
1743                if stream.expect_string(".m64n192k32").is_ok() {
1744                    return Ok(Shape::M64n192k32);
1745                }
1746                stream.set_position(saved_pos);
1747            }
1748            stream.set_position(saved_pos);
1749            let saved_pos = stream.position();
1750            // Try M64n200k32
1751            {
1752                let saved_pos = stream.position();
1753                if stream.expect_string(".m64n200k32").is_ok() {
1754                    return Ok(Shape::M64n200k32);
1755                }
1756                stream.set_position(saved_pos);
1757            }
1758            stream.set_position(saved_pos);
1759            let saved_pos = stream.position();
1760            // Try M64n208k32
1761            {
1762                let saved_pos = stream.position();
1763                if stream.expect_string(".m64n208k32").is_ok() {
1764                    return Ok(Shape::M64n208k32);
1765                }
1766                stream.set_position(saved_pos);
1767            }
1768            stream.set_position(saved_pos);
1769            let saved_pos = stream.position();
1770            // Try M64n216k32
1771            {
1772                let saved_pos = stream.position();
1773                if stream.expect_string(".m64n216k32").is_ok() {
1774                    return Ok(Shape::M64n216k32);
1775                }
1776                stream.set_position(saved_pos);
1777            }
1778            stream.set_position(saved_pos);
1779            let saved_pos = stream.position();
1780            // Try M64n224k32
1781            {
1782                let saved_pos = stream.position();
1783                if stream.expect_string(".m64n224k32").is_ok() {
1784                    return Ok(Shape::M64n224k32);
1785                }
1786                stream.set_position(saved_pos);
1787            }
1788            stream.set_position(saved_pos);
1789            let saved_pos = stream.position();
1790            // Try M64n232k32
1791            {
1792                let saved_pos = stream.position();
1793                if stream.expect_string(".m64n232k32").is_ok() {
1794                    return Ok(Shape::M64n232k32);
1795                }
1796                stream.set_position(saved_pos);
1797            }
1798            stream.set_position(saved_pos);
1799            let saved_pos = stream.position();
1800            // Try M64n240k32
1801            {
1802                let saved_pos = stream.position();
1803                if stream.expect_string(".m64n240k32").is_ok() {
1804                    return Ok(Shape::M64n240k32);
1805                }
1806                stream.set_position(saved_pos);
1807            }
1808            stream.set_position(saved_pos);
1809            let saved_pos = stream.position();
1810            // Try M64n248k32
1811            {
1812                let saved_pos = stream.position();
1813                if stream.expect_string(".m64n248k32").is_ok() {
1814                    return Ok(Shape::M64n248k32);
1815                }
1816                stream.set_position(saved_pos);
1817            }
1818            stream.set_position(saved_pos);
1819            let saved_pos = stream.position();
1820            // Try M64n256k32
1821            {
1822                let saved_pos = stream.position();
1823                if stream.expect_string(".m64n256k32").is_ok() {
1824                    return Ok(Shape::M64n256k32);
1825                }
1826                stream.set_position(saved_pos);
1827            }
1828            stream.set_position(saved_pos);
1829            let saved_pos = stream.position();
1830            // Try M64n16k32
1831            {
1832                let saved_pos = stream.position();
1833                if stream.expect_string(".m64n16k32").is_ok() {
1834                    return Ok(Shape::M64n16k32);
1835                }
1836                stream.set_position(saved_pos);
1837            }
1838            stream.set_position(saved_pos);
1839            let saved_pos = stream.position();
1840            // Try M64n24k32
1841            {
1842                let saved_pos = stream.position();
1843                if stream.expect_string(".m64n24k32").is_ok() {
1844                    return Ok(Shape::M64n24k32);
1845                }
1846                stream.set_position(saved_pos);
1847            }
1848            stream.set_position(saved_pos);
1849            let saved_pos = stream.position();
1850            // Try M64n32k32
1851            {
1852                let saved_pos = stream.position();
1853                if stream.expect_string(".m64n32k32").is_ok() {
1854                    return Ok(Shape::M64n32k32);
1855                }
1856                stream.set_position(saved_pos);
1857            }
1858            stream.set_position(saved_pos);
1859            let saved_pos = stream.position();
1860            // Try M64n40k32
1861            {
1862                let saved_pos = stream.position();
1863                if stream.expect_string(".m64n40k32").is_ok() {
1864                    return Ok(Shape::M64n40k32);
1865                }
1866                stream.set_position(saved_pos);
1867            }
1868            stream.set_position(saved_pos);
1869            let saved_pos = stream.position();
1870            // Try M64n48k32
1871            {
1872                let saved_pos = stream.position();
1873                if stream.expect_string(".m64n48k32").is_ok() {
1874                    return Ok(Shape::M64n48k32);
1875                }
1876                stream.set_position(saved_pos);
1877            }
1878            stream.set_position(saved_pos);
1879            let saved_pos = stream.position();
1880            // Try M64n56k32
1881            {
1882                let saved_pos = stream.position();
1883                if stream.expect_string(".m64n56k32").is_ok() {
1884                    return Ok(Shape::M64n56k32);
1885                }
1886                stream.set_position(saved_pos);
1887            }
1888            stream.set_position(saved_pos);
1889            let saved_pos = stream.position();
1890            // Try M64n64k32
1891            {
1892                let saved_pos = stream.position();
1893                if stream.expect_string(".m64n64k32").is_ok() {
1894                    return Ok(Shape::M64n64k32);
1895                }
1896                stream.set_position(saved_pos);
1897            }
1898            stream.set_position(saved_pos);
1899            let saved_pos = stream.position();
1900            // Try M64n72k32
1901            {
1902                let saved_pos = stream.position();
1903                if stream.expect_string(".m64n72k32").is_ok() {
1904                    return Ok(Shape::M64n72k32);
1905                }
1906                stream.set_position(saved_pos);
1907            }
1908            stream.set_position(saved_pos);
1909            let saved_pos = stream.position();
1910            // Try M64n80k32
1911            {
1912                let saved_pos = stream.position();
1913                if stream.expect_string(".m64n80k32").is_ok() {
1914                    return Ok(Shape::M64n80k32);
1915                }
1916                stream.set_position(saved_pos);
1917            }
1918            stream.set_position(saved_pos);
1919            let saved_pos = stream.position();
1920            // Try M64n88k32
1921            {
1922                let saved_pos = stream.position();
1923                if stream.expect_string(".m64n88k32").is_ok() {
1924                    return Ok(Shape::M64n88k32);
1925                }
1926                stream.set_position(saved_pos);
1927            }
1928            stream.set_position(saved_pos);
1929            let saved_pos = stream.position();
1930            // Try M64n96k32
1931            {
1932                let saved_pos = stream.position();
1933                if stream.expect_string(".m64n96k32").is_ok() {
1934                    return Ok(Shape::M64n96k32);
1935                }
1936                stream.set_position(saved_pos);
1937            }
1938            stream.set_position(saved_pos);
1939            let saved_pos = stream.position();
1940            // Try M64n8k32
1941            {
1942                let saved_pos = stream.position();
1943                if stream.expect_string(".m64n8k32").is_ok() {
1944                    return Ok(Shape::M64n8k32);
1945                }
1946                stream.set_position(saved_pos);
1947            }
1948            stream.set_position(saved_pos);
1949            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1950            let expected = &[".m64n104k32", ".m64n112k32", ".m64n120k32", ".m64n128k32", ".m64n136k32", ".m64n144k32", ".m64n152k32", ".m64n160k32", ".m64n168k32", ".m64n176k32", ".m64n184k32", ".m64n192k32", ".m64n200k32", ".m64n208k32", ".m64n216k32", ".m64n224k32", ".m64n232k32", ".m64n240k32", ".m64n248k32", ".m64n256k32", ".m64n16k32", ".m64n24k32", ".m64n32k32", ".m64n40k32", ".m64n48k32", ".m64n56k32", ".m64n64k32", ".m64n72k32", ".m64n80k32", ".m64n88k32", ".m64n96k32", ".m64n8k32"];
1951            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1952            Err(crate::parser::unexpected_value(span, expected, found))
1953        }
1954    }
1955
1956    impl PtxParser for WgmmaMmaAsyncSyncAlignedShapeDtypeAtypeBtype {
1957        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1958            stream.expect_string("wgmma")?;
1959            stream.expect_string(".mma_async")?;
1960            let mma_async = ();
1961            stream.expect_complete()?;
1962            stream.expect_string(".sync")?;
1963            let sync = ();
1964            stream.expect_complete()?;
1965            stream.expect_string(".aligned")?;
1966            let aligned = ();
1967            stream.expect_complete()?;
1968            let shape = Shape::parse(stream)?;
1969            stream.expect_complete()?;
1970            let dtype = Dtype::parse(stream)?;
1971            stream.expect_complete()?;
1972            let atype = Atype::parse(stream)?;
1973            stream.expect_complete()?;
1974            let btype = Btype::parse(stream)?;
1975            stream.expect_complete()?;
1976            let d = GeneralOperand::parse(stream)?;
1977            stream.expect_complete()?;
1978            stream.expect(&PtxToken::Comma)?;
1979            let a_desc = GeneralOperand::parse(stream)?;
1980            stream.expect_complete()?;
1981            stream.expect(&PtxToken::Comma)?;
1982            let b_desc = GeneralOperand::parse(stream)?;
1983            stream.expect_complete()?;
1984            stream.expect(&PtxToken::Comma)?;
1985            let scale_d = GeneralOperand::parse(stream)?;
1986            stream.expect_complete()?;
1987            stream.expect(&PtxToken::Comma)?;
1988            let imm_scale_a = GeneralOperand::parse(stream)?;
1989            stream.expect_complete()?;
1990            stream.expect(&PtxToken::Comma)?;
1991            let imm_scale_b = GeneralOperand::parse(stream)?;
1992            stream.expect_complete()?;
1993            stream.expect_complete()?;
1994            stream.expect(&PtxToken::Semicolon)?;
1995            Ok(WgmmaMmaAsyncSyncAlignedShapeDtypeAtypeBtype {
1996                mma_async,
1997                sync,
1998                aligned,
1999                shape,
2000                dtype,
2001                atype,
2002                btype,
2003                d,
2004                a_desc,
2005                b_desc,
2006                scale_d,
2007                imm_scale_a,
2008                imm_scale_b,
2009            })
2010        }
2011    }
2012
2013
2014    impl PtxParser for WgmmaMmaAsyncSyncAlignedShapeDtypeAtypeBtype1 {
2015        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2016            stream.expect_string("wgmma")?;
2017            stream.expect_string(".mma_async")?;
2018            let mma_async = ();
2019            stream.expect_complete()?;
2020            stream.expect_string(".sync")?;
2021            let sync = ();
2022            stream.expect_complete()?;
2023            stream.expect_string(".aligned")?;
2024            let aligned = ();
2025            stream.expect_complete()?;
2026            let shape = Shape::parse(stream)?;
2027            stream.expect_complete()?;
2028            let dtype = Dtype::parse(stream)?;
2029            stream.expect_complete()?;
2030            let atype = Atype::parse(stream)?;
2031            stream.expect_complete()?;
2032            let btype = Btype::parse(stream)?;
2033            stream.expect_complete()?;
2034            let d = GeneralOperand::parse(stream)?;
2035            stream.expect_complete()?;
2036            stream.expect(&PtxToken::Comma)?;
2037            let a = GeneralOperand::parse(stream)?;
2038            stream.expect_complete()?;
2039            stream.expect(&PtxToken::Comma)?;
2040            let b_desc = GeneralOperand::parse(stream)?;
2041            stream.expect_complete()?;
2042            stream.expect(&PtxToken::Comma)?;
2043            let scale_d = GeneralOperand::parse(stream)?;
2044            stream.expect_complete()?;
2045            stream.expect(&PtxToken::Comma)?;
2046            let imm_scale_a = GeneralOperand::parse(stream)?;
2047            stream.expect_complete()?;
2048            stream.expect(&PtxToken::Comma)?;
2049            let imm_scale_b = GeneralOperand::parse(stream)?;
2050            stream.expect_complete()?;
2051            stream.expect_complete()?;
2052            stream.expect(&PtxToken::Semicolon)?;
2053            Ok(WgmmaMmaAsyncSyncAlignedShapeDtypeAtypeBtype1 {
2054                mma_async,
2055                sync,
2056                aligned,
2057                shape,
2058                dtype,
2059                atype,
2060                btype,
2061                d,
2062                a,
2063                b_desc,
2064                scale_d,
2065                imm_scale_a,
2066                imm_scale_b,
2067            })
2068        }
2069    }
2070
2071
2072}
2073
2074pub mod section_4 {
2075    use super::*;
2076    use crate::r#type::instruction::wgmma_mma_async::section_4::*;
2077
2078    // ============================================================================
2079    // Generated enum parsers
2080    // ============================================================================
2081
2082    impl PtxParser for Atype {
2083        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2084            // Try S8
2085            {
2086                let saved_pos = stream.position();
2087                if stream.expect_string(".s8").is_ok() {
2088                    return Ok(Atype::S8);
2089                }
2090                stream.set_position(saved_pos);
2091            }
2092            let saved_pos = stream.position();
2093            // Try U8
2094            {
2095                let saved_pos = stream.position();
2096                if stream.expect_string(".u8").is_ok() {
2097                    return Ok(Atype::U8);
2098                }
2099                stream.set_position(saved_pos);
2100            }
2101            stream.set_position(saved_pos);
2102            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
2103            let expected = &[".s8", ".u8"];
2104            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
2105            Err(crate::parser::unexpected_value(span, expected, found))
2106        }
2107    }
2108
2109    impl PtxParser for Btype {
2110        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2111            // Try S8
2112            {
2113                let saved_pos = stream.position();
2114                if stream.expect_string(".s8").is_ok() {
2115                    return Ok(Btype::S8);
2116                }
2117                stream.set_position(saved_pos);
2118            }
2119            let saved_pos = stream.position();
2120            // Try U8
2121            {
2122                let saved_pos = stream.position();
2123                if stream.expect_string(".u8").is_ok() {
2124                    return Ok(Btype::U8);
2125                }
2126                stream.set_position(saved_pos);
2127            }
2128            stream.set_position(saved_pos);
2129            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
2130            let expected = &[".s8", ".u8"];
2131            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
2132            Err(crate::parser::unexpected_value(span, expected, found))
2133        }
2134    }
2135
2136    impl PtxParser for Shape {
2137        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2138            // Try M64n112k32
2139            {
2140                let saved_pos = stream.position();
2141                if stream.expect_string(".m64n112k32").is_ok() {
2142                    return Ok(Shape::M64n112k32);
2143                }
2144                stream.set_position(saved_pos);
2145            }
2146            let saved_pos = stream.position();
2147            // Try M64n128k32
2148            {
2149                let saved_pos = stream.position();
2150                if stream.expect_string(".m64n128k32").is_ok() {
2151                    return Ok(Shape::M64n128k32);
2152                }
2153                stream.set_position(saved_pos);
2154            }
2155            stream.set_position(saved_pos);
2156            let saved_pos = stream.position();
2157            // Try M64n144k32
2158            {
2159                let saved_pos = stream.position();
2160                if stream.expect_string(".m64n144k32").is_ok() {
2161                    return Ok(Shape::M64n144k32);
2162                }
2163                stream.set_position(saved_pos);
2164            }
2165            stream.set_position(saved_pos);
2166            let saved_pos = stream.position();
2167            // Try M64n160k32
2168            {
2169                let saved_pos = stream.position();
2170                if stream.expect_string(".m64n160k32").is_ok() {
2171                    return Ok(Shape::M64n160k32);
2172                }
2173                stream.set_position(saved_pos);
2174            }
2175            stream.set_position(saved_pos);
2176            let saved_pos = stream.position();
2177            // Try M64n176k32
2178            {
2179                let saved_pos = stream.position();
2180                if stream.expect_string(".m64n176k32").is_ok() {
2181                    return Ok(Shape::M64n176k32);
2182                }
2183                stream.set_position(saved_pos);
2184            }
2185            stream.set_position(saved_pos);
2186            let saved_pos = stream.position();
2187            // Try M64n192k32
2188            {
2189                let saved_pos = stream.position();
2190                if stream.expect_string(".m64n192k32").is_ok() {
2191                    return Ok(Shape::M64n192k32);
2192                }
2193                stream.set_position(saved_pos);
2194            }
2195            stream.set_position(saved_pos);
2196            let saved_pos = stream.position();
2197            // Try M64n208k32
2198            {
2199                let saved_pos = stream.position();
2200                if stream.expect_string(".m64n208k32").is_ok() {
2201                    return Ok(Shape::M64n208k32);
2202                }
2203                stream.set_position(saved_pos);
2204            }
2205            stream.set_position(saved_pos);
2206            let saved_pos = stream.position();
2207            // Try M64n224k32
2208            {
2209                let saved_pos = stream.position();
2210                if stream.expect_string(".m64n224k32").is_ok() {
2211                    return Ok(Shape::M64n224k32);
2212                }
2213                stream.set_position(saved_pos);
2214            }
2215            stream.set_position(saved_pos);
2216            let saved_pos = stream.position();
2217            // Try M64n16k32
2218            {
2219                let saved_pos = stream.position();
2220                if stream.expect_string(".m64n16k32").is_ok() {
2221                    return Ok(Shape::M64n16k32);
2222                }
2223                stream.set_position(saved_pos);
2224            }
2225            stream.set_position(saved_pos);
2226            let saved_pos = stream.position();
2227            // Try M64n24k32
2228            {
2229                let saved_pos = stream.position();
2230                if stream.expect_string(".m64n24k32").is_ok() {
2231                    return Ok(Shape::M64n24k32);
2232                }
2233                stream.set_position(saved_pos);
2234            }
2235            stream.set_position(saved_pos);
2236            let saved_pos = stream.position();
2237            // Try M64n32k32
2238            {
2239                let saved_pos = stream.position();
2240                if stream.expect_string(".m64n32k32").is_ok() {
2241                    return Ok(Shape::M64n32k32);
2242                }
2243                stream.set_position(saved_pos);
2244            }
2245            stream.set_position(saved_pos);
2246            let saved_pos = stream.position();
2247            // Try M64n48k32
2248            {
2249                let saved_pos = stream.position();
2250                if stream.expect_string(".m64n48k32").is_ok() {
2251                    return Ok(Shape::M64n48k32);
2252                }
2253                stream.set_position(saved_pos);
2254            }
2255            stream.set_position(saved_pos);
2256            let saved_pos = stream.position();
2257            // Try M64n64k32
2258            {
2259                let saved_pos = stream.position();
2260                if stream.expect_string(".m64n64k32").is_ok() {
2261                    return Ok(Shape::M64n64k32);
2262                }
2263                stream.set_position(saved_pos);
2264            }
2265            stream.set_position(saved_pos);
2266            let saved_pos = stream.position();
2267            // Try M64n80k32
2268            {
2269                let saved_pos = stream.position();
2270                if stream.expect_string(".m64n80k32").is_ok() {
2271                    return Ok(Shape::M64n80k32);
2272                }
2273                stream.set_position(saved_pos);
2274            }
2275            stream.set_position(saved_pos);
2276            let saved_pos = stream.position();
2277            // Try M64n96k32
2278            {
2279                let saved_pos = stream.position();
2280                if stream.expect_string(".m64n96k32").is_ok() {
2281                    return Ok(Shape::M64n96k32);
2282                }
2283                stream.set_position(saved_pos);
2284            }
2285            stream.set_position(saved_pos);
2286            let saved_pos = stream.position();
2287            // Try M64n8k32
2288            {
2289                let saved_pos = stream.position();
2290                if stream.expect_string(".m64n8k32").is_ok() {
2291                    return Ok(Shape::M64n8k32);
2292                }
2293                stream.set_position(saved_pos);
2294            }
2295            stream.set_position(saved_pos);
2296            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
2297            let expected = &[".m64n112k32", ".m64n128k32", ".m64n144k32", ".m64n160k32", ".m64n176k32", ".m64n192k32", ".m64n208k32", ".m64n224k32", ".m64n16k32", ".m64n24k32", ".m64n32k32", ".m64n48k32", ".m64n64k32", ".m64n80k32", ".m64n96k32", ".m64n8k32"];
2298            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
2299            Err(crate::parser::unexpected_value(span, expected, found))
2300        }
2301    }
2302
2303    impl PtxParser for WgmmaMmaAsyncSyncAlignedShapeSatfiniteS32AtypeBtype {
2304        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2305            stream.expect_string("wgmma")?;
2306            stream.expect_string(".mma_async")?;
2307            let mma_async = ();
2308            stream.expect_complete()?;
2309            stream.expect_string(".sync")?;
2310            let sync = ();
2311            stream.expect_complete()?;
2312            stream.expect_string(".aligned")?;
2313            let aligned = ();
2314            stream.expect_complete()?;
2315            let shape = Shape::parse(stream)?;
2316            stream.expect_complete()?;
2317            let saved_pos = stream.position();
2318            let satfinite = stream.expect_string(".satfinite").is_ok();
2319            if !satfinite {
2320                stream.set_position(saved_pos);
2321            }
2322            stream.expect_complete()?;
2323            stream.expect_string(".s32")?;
2324            let s32 = ();
2325            stream.expect_complete()?;
2326            let atype = Atype::parse(stream)?;
2327            stream.expect_complete()?;
2328            let btype = Btype::parse(stream)?;
2329            stream.expect_complete()?;
2330            let d = GeneralOperand::parse(stream)?;
2331            stream.expect_complete()?;
2332            stream.expect(&PtxToken::Comma)?;
2333            let a_desc = GeneralOperand::parse(stream)?;
2334            stream.expect_complete()?;
2335            stream.expect(&PtxToken::Comma)?;
2336            let b_desc = GeneralOperand::parse(stream)?;
2337            stream.expect_complete()?;
2338            stream.expect(&PtxToken::Comma)?;
2339            let scale_d = GeneralOperand::parse(stream)?;
2340            stream.expect_complete()?;
2341            stream.expect_complete()?;
2342            stream.expect(&PtxToken::Semicolon)?;
2343            Ok(WgmmaMmaAsyncSyncAlignedShapeSatfiniteS32AtypeBtype {
2344                mma_async,
2345                sync,
2346                aligned,
2347                shape,
2348                satfinite,
2349                s32,
2350                atype,
2351                btype,
2352                d,
2353                a_desc,
2354                b_desc,
2355                scale_d,
2356            })
2357        }
2358    }
2359
2360
2361    impl PtxParser for WgmmaMmaAsyncSyncAlignedShapeSatfiniteS32AtypeBtype1 {
2362        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2363            stream.expect_string("wgmma")?;
2364            stream.expect_string(".mma_async")?;
2365            let mma_async = ();
2366            stream.expect_complete()?;
2367            stream.expect_string(".sync")?;
2368            let sync = ();
2369            stream.expect_complete()?;
2370            stream.expect_string(".aligned")?;
2371            let aligned = ();
2372            stream.expect_complete()?;
2373            let shape = Shape::parse(stream)?;
2374            stream.expect_complete()?;
2375            let saved_pos = stream.position();
2376            let satfinite = stream.expect_string(".satfinite").is_ok();
2377            if !satfinite {
2378                stream.set_position(saved_pos);
2379            }
2380            stream.expect_complete()?;
2381            stream.expect_string(".s32")?;
2382            let s32 = ();
2383            stream.expect_complete()?;
2384            let atype = Atype::parse(stream)?;
2385            stream.expect_complete()?;
2386            let btype = Btype::parse(stream)?;
2387            stream.expect_complete()?;
2388            let d = GeneralOperand::parse(stream)?;
2389            stream.expect_complete()?;
2390            stream.expect(&PtxToken::Comma)?;
2391            let a = GeneralOperand::parse(stream)?;
2392            stream.expect_complete()?;
2393            stream.expect(&PtxToken::Comma)?;
2394            let b_desc = GeneralOperand::parse(stream)?;
2395            stream.expect_complete()?;
2396            stream.expect(&PtxToken::Comma)?;
2397            let scale_d = GeneralOperand::parse(stream)?;
2398            stream.expect_complete()?;
2399            stream.expect_complete()?;
2400            stream.expect(&PtxToken::Semicolon)?;
2401            Ok(WgmmaMmaAsyncSyncAlignedShapeSatfiniteS32AtypeBtype1 {
2402                mma_async,
2403                sync,
2404                aligned,
2405                shape,
2406                satfinite,
2407                s32,
2408                atype,
2409                btype,
2410                d,
2411                a,
2412                b_desc,
2413                scale_d,
2414            })
2415        }
2416    }
2417
2418
2419}
2420
2421pub mod section_5 {
2422    use super::*;
2423    use crate::r#type::instruction::wgmma_mma_async::section_5::*;
2424
2425    // ============================================================================
2426    // Generated enum parsers
2427    // ============================================================================
2428
2429    impl PtxParser for Op {
2430        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2431            // Try And
2432            {
2433                let saved_pos = stream.position();
2434                if stream.expect_string(".and").is_ok() {
2435                    return Ok(Op::And);
2436                }
2437                stream.set_position(saved_pos);
2438            }
2439            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
2440            let expected = &[".and"];
2441            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
2442            Err(crate::parser::unexpected_value(span, expected, found))
2443        }
2444    }
2445
2446    impl PtxParser for Shape {
2447        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2448            // Try M64n112k256
2449            {
2450                let saved_pos = stream.position();
2451                if stream.expect_string(".m64n112k256").is_ok() {
2452                    return Ok(Shape::M64n112k256);
2453                }
2454                stream.set_position(saved_pos);
2455            }
2456            let saved_pos = stream.position();
2457            // Try M64n128k256
2458            {
2459                let saved_pos = stream.position();
2460                if stream.expect_string(".m64n128k256").is_ok() {
2461                    return Ok(Shape::M64n128k256);
2462                }
2463                stream.set_position(saved_pos);
2464            }
2465            stream.set_position(saved_pos);
2466            let saved_pos = stream.position();
2467            // Try M64n144k256
2468            {
2469                let saved_pos = stream.position();
2470                if stream.expect_string(".m64n144k256").is_ok() {
2471                    return Ok(Shape::M64n144k256);
2472                }
2473                stream.set_position(saved_pos);
2474            }
2475            stream.set_position(saved_pos);
2476            let saved_pos = stream.position();
2477            // Try M64n160k256
2478            {
2479                let saved_pos = stream.position();
2480                if stream.expect_string(".m64n160k256").is_ok() {
2481                    return Ok(Shape::M64n160k256);
2482                }
2483                stream.set_position(saved_pos);
2484            }
2485            stream.set_position(saved_pos);
2486            let saved_pos = stream.position();
2487            // Try M64n176k256
2488            {
2489                let saved_pos = stream.position();
2490                if stream.expect_string(".m64n176k256").is_ok() {
2491                    return Ok(Shape::M64n176k256);
2492                }
2493                stream.set_position(saved_pos);
2494            }
2495            stream.set_position(saved_pos);
2496            let saved_pos = stream.position();
2497            // Try M64n192k256
2498            {
2499                let saved_pos = stream.position();
2500                if stream.expect_string(".m64n192k256").is_ok() {
2501                    return Ok(Shape::M64n192k256);
2502                }
2503                stream.set_position(saved_pos);
2504            }
2505            stream.set_position(saved_pos);
2506            let saved_pos = stream.position();
2507            // Try M64n208k256
2508            {
2509                let saved_pos = stream.position();
2510                if stream.expect_string(".m64n208k256").is_ok() {
2511                    return Ok(Shape::M64n208k256);
2512                }
2513                stream.set_position(saved_pos);
2514            }
2515            stream.set_position(saved_pos);
2516            let saved_pos = stream.position();
2517            // Try M64n224k256
2518            {
2519                let saved_pos = stream.position();
2520                if stream.expect_string(".m64n224k256").is_ok() {
2521                    return Ok(Shape::M64n224k256);
2522                }
2523                stream.set_position(saved_pos);
2524            }
2525            stream.set_position(saved_pos);
2526            let saved_pos = stream.position();
2527            // Try M64n240k256
2528            {
2529                let saved_pos = stream.position();
2530                if stream.expect_string(".m64n240k256").is_ok() {
2531                    return Ok(Shape::M64n240k256);
2532                }
2533                stream.set_position(saved_pos);
2534            }
2535            stream.set_position(saved_pos);
2536            let saved_pos = stream.position();
2537            // Try M64n256k256
2538            {
2539                let saved_pos = stream.position();
2540                if stream.expect_string(".m64n256k256").is_ok() {
2541                    return Ok(Shape::M64n256k256);
2542                }
2543                stream.set_position(saved_pos);
2544            }
2545            stream.set_position(saved_pos);
2546            let saved_pos = stream.position();
2547            // Try M64n16k256
2548            {
2549                let saved_pos = stream.position();
2550                if stream.expect_string(".m64n16k256").is_ok() {
2551                    return Ok(Shape::M64n16k256);
2552                }
2553                stream.set_position(saved_pos);
2554            }
2555            stream.set_position(saved_pos);
2556            let saved_pos = stream.position();
2557            // Try M64n24k256
2558            {
2559                let saved_pos = stream.position();
2560                if stream.expect_string(".m64n24k256").is_ok() {
2561                    return Ok(Shape::M64n24k256);
2562                }
2563                stream.set_position(saved_pos);
2564            }
2565            stream.set_position(saved_pos);
2566            let saved_pos = stream.position();
2567            // Try M64n32k256
2568            {
2569                let saved_pos = stream.position();
2570                if stream.expect_string(".m64n32k256").is_ok() {
2571                    return Ok(Shape::M64n32k256);
2572                }
2573                stream.set_position(saved_pos);
2574            }
2575            stream.set_position(saved_pos);
2576            let saved_pos = stream.position();
2577            // Try M64n48k256
2578            {
2579                let saved_pos = stream.position();
2580                if stream.expect_string(".m64n48k256").is_ok() {
2581                    return Ok(Shape::M64n48k256);
2582                }
2583                stream.set_position(saved_pos);
2584            }
2585            stream.set_position(saved_pos);
2586            let saved_pos = stream.position();
2587            // Try M64n64k256
2588            {
2589                let saved_pos = stream.position();
2590                if stream.expect_string(".m64n64k256").is_ok() {
2591                    return Ok(Shape::M64n64k256);
2592                }
2593                stream.set_position(saved_pos);
2594            }
2595            stream.set_position(saved_pos);
2596            let saved_pos = stream.position();
2597            // Try M64n80k256
2598            {
2599                let saved_pos = stream.position();
2600                if stream.expect_string(".m64n80k256").is_ok() {
2601                    return Ok(Shape::M64n80k256);
2602                }
2603                stream.set_position(saved_pos);
2604            }
2605            stream.set_position(saved_pos);
2606            let saved_pos = stream.position();
2607            // Try M64n96k256
2608            {
2609                let saved_pos = stream.position();
2610                if stream.expect_string(".m64n96k256").is_ok() {
2611                    return Ok(Shape::M64n96k256);
2612                }
2613                stream.set_position(saved_pos);
2614            }
2615            stream.set_position(saved_pos);
2616            let saved_pos = stream.position();
2617            // Try M64n8k256
2618            {
2619                let saved_pos = stream.position();
2620                if stream.expect_string(".m64n8k256").is_ok() {
2621                    return Ok(Shape::M64n8k256);
2622                }
2623                stream.set_position(saved_pos);
2624            }
2625            stream.set_position(saved_pos);
2626            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
2627            let expected = &[".m64n112k256", ".m64n128k256", ".m64n144k256", ".m64n160k256", ".m64n176k256", ".m64n192k256", ".m64n208k256", ".m64n224k256", ".m64n240k256", ".m64n256k256", ".m64n16k256", ".m64n24k256", ".m64n32k256", ".m64n48k256", ".m64n64k256", ".m64n80k256", ".m64n96k256", ".m64n8k256"];
2628            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
2629            Err(crate::parser::unexpected_value(span, expected, found))
2630        }
2631    }
2632
2633    impl PtxParser for WgmmaMmaAsyncSyncAlignedShapeS32B1B1OpPopc {
2634        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2635            stream.expect_string("wgmma")?;
2636            stream.expect_string(".mma_async")?;
2637            let mma_async = ();
2638            stream.expect_complete()?;
2639            stream.expect_string(".sync")?;
2640            let sync = ();
2641            stream.expect_complete()?;
2642            stream.expect_string(".aligned")?;
2643            let aligned = ();
2644            stream.expect_complete()?;
2645            let shape = Shape::parse(stream)?;
2646            stream.expect_complete()?;
2647            stream.expect_string(".s32")?;
2648            let s32 = ();
2649            stream.expect_complete()?;
2650            stream.expect_string(".b1")?;
2651            let b1 = ();
2652            stream.expect_complete()?;
2653            stream.expect_string(".b1")?;
2654            let b12 = ();
2655            stream.expect_complete()?;
2656            let op = Op::parse(stream)?;
2657            stream.expect_complete()?;
2658            stream.expect_string(".popc")?;
2659            let popc = ();
2660            stream.expect_complete()?;
2661            let d = GeneralOperand::parse(stream)?;
2662            stream.expect_complete()?;
2663            stream.expect(&PtxToken::Comma)?;
2664            let a_desc = GeneralOperand::parse(stream)?;
2665            stream.expect_complete()?;
2666            stream.expect(&PtxToken::Comma)?;
2667            let b_desc = GeneralOperand::parse(stream)?;
2668            stream.expect_complete()?;
2669            stream.expect(&PtxToken::Comma)?;
2670            let scale_d = GeneralOperand::parse(stream)?;
2671            stream.expect_complete()?;
2672            stream.expect_complete()?;
2673            stream.expect(&PtxToken::Semicolon)?;
2674            Ok(WgmmaMmaAsyncSyncAlignedShapeS32B1B1OpPopc {
2675                mma_async,
2676                sync,
2677                aligned,
2678                shape,
2679                s32,
2680                b1,
2681                b12,
2682                op,
2683                popc,
2684                d,
2685                a_desc,
2686                b_desc,
2687                scale_d,
2688            })
2689        }
2690    }
2691
2692
2693    impl PtxParser for WgmmaMmaAsyncSyncAlignedShapeS32B1B1OpPopc1 {
2694        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2695            stream.expect_string("wgmma")?;
2696            stream.expect_string(".mma_async")?;
2697            let mma_async = ();
2698            stream.expect_complete()?;
2699            stream.expect_string(".sync")?;
2700            let sync = ();
2701            stream.expect_complete()?;
2702            stream.expect_string(".aligned")?;
2703            let aligned = ();
2704            stream.expect_complete()?;
2705            let shape = Shape::parse(stream)?;
2706            stream.expect_complete()?;
2707            stream.expect_string(".s32")?;
2708            let s32 = ();
2709            stream.expect_complete()?;
2710            stream.expect_string(".b1")?;
2711            let b1 = ();
2712            stream.expect_complete()?;
2713            stream.expect_string(".b1")?;
2714            let b12 = ();
2715            stream.expect_complete()?;
2716            let op = Op::parse(stream)?;
2717            stream.expect_complete()?;
2718            stream.expect_string(".popc")?;
2719            let popc = ();
2720            stream.expect_complete()?;
2721            let d = GeneralOperand::parse(stream)?;
2722            stream.expect_complete()?;
2723            stream.expect(&PtxToken::Comma)?;
2724            let a = GeneralOperand::parse(stream)?;
2725            stream.expect_complete()?;
2726            stream.expect(&PtxToken::Comma)?;
2727            let b_desc = GeneralOperand::parse(stream)?;
2728            stream.expect_complete()?;
2729            stream.expect(&PtxToken::Comma)?;
2730            let scale_d = GeneralOperand::parse(stream)?;
2731            stream.expect_complete()?;
2732            stream.expect_complete()?;
2733            stream.expect(&PtxToken::Semicolon)?;
2734            Ok(WgmmaMmaAsyncSyncAlignedShapeS32B1B1OpPopc1 {
2735                mma_async,
2736                sync,
2737                aligned,
2738                shape,
2739                s32,
2740                b1,
2741                b12,
2742                op,
2743                popc,
2744                d,
2745                a,
2746                b_desc,
2747                scale_d,
2748            })
2749        }
2750    }
2751
2752
2753}
2754
ptx_parser/parser/instruction/wgmma_mma_async.rs

ptx_parser/parser/instruction/
wgmma_mma_async.rs
