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

ptx_parser/parser/instruction/
wgmma_mma_async.rs
