ptx_parser/parser/instruction/
wgmma_mma_async_sp.rs

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

ptx_parser/parser/instruction/
wgmma_mma_async_sp.rs
