ptx_parser/parser/instruction/
wmma_mma.rs

1//! Original PTX specification:
2//!
3//! // Floating point (.f16 multiplicands) wmma.mma
4//! wmma.mma.sync.aligned.alayout.blayout.shape.dtype.ctype d, a, b, c;
5//! ----------------------------------------------------------------
6//! // Integer (.u8/.s8 multiplicands) wmma.mma
7//! wmma.mma.sync.aligned.alayout.blayout.shape.s32.atype.btype.s32{.satfinite} d, a, b, c;
8//! .alayout = {.row, .col};
9//! .blayout = {.row, .col};
10//! .shape  =  {.m16n16k16, .m8n32k16, .m32n8k16};
11//! .dtype   = {.f16, .f32};
12//! .atype   = {.s8, .u8};
13//! .btype   = {.s8, .u8};
14//! .ctype   = {.f16, .f32};
15//! ----------------------------------------------------------------
16//! // Floating point format .bf16 wmma.mma:
17//! wmma.mma.sync.aligned.alayout.blayout.shape.f32.atype.btype.f32 d, a, b, c;
18//! .alayout = {.row, .col};
19//! .blayout = {.row, .col};
20//! .shape   = {.m16n16k16, .m8n32k16, .m32n8k16};
21//! .atype   = {.bf16 };
22//! .btype   = {.bf16};
23//! ----------------------------------------------------------------
24//! // Floating point format .tf32 wmma.mma:
25//! wmma.mma.sync.aligned.alayout.blayout.shape.f32.atype.btype.f32 d, a, b, c;
26//! .alayout = {.row, .col};
27//! .blayout = {.row, .col};
28//! .shape   = {.m16n16k8 };
29//! .atype   = {.tf32 };
30//! .btype   = {.tf32};
31//! ----------------------------------------------------------------
32//! // Floating point Double precision wmma.mma:
33//! wmma.mma.sync.aligned.alayout.blayout.shape{.rnd}.f64.f64.f64.f64 d, a, b, c;
34//! .alayout = {.row, .col};
35//! .blayout = {.row, .col};
36//! .shape   = {.m8n8k4 };
37//! .rnd = { .rn, .rz, .rm, .rp };
38//! ----------------------------------------------------------------
39//! // Sub-byte (.u4/.s4 multiplicands) wmma.mma:
40//! wmma.mma.sync.aligned.row.col.shape.s32.atype.btype.s32{.satfinite} d, a, b, c;
41//! .shape  = {.m8n8k32};
42//! .atype  = {.s4, .u4};
43//! .btype  = {.s4, .u4};
44//! ----------------------------------------------------------------
45//! // Single-bit (.b1 multiplicands) wmma.mma:
46//! wmma.mma.op.popc.sync.aligned.row.col.shape.s32.atype.btype.s32 d, a, b, c;
47//! .shape  = {.m8n8k128};
48//! .atype  = {.b1};
49//! .btype  = {.b1};
50//! .op     = {.xor, .and};
51
52#![allow(unused)]
53
54use crate::lexer::PtxToken;
55use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
56use crate::r#type::common::*;
57
58pub mod section_0 {
59    use super::*;
60    use crate::r#type::instruction::wmma_mma::section_0::*;
61
62    impl PtxParser for WmmaMmaSyncAlignedAlayoutBlayoutShapeDtypeCtype {
63        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
64            stream.expect_string("wmma")?;
65            stream.expect_string(".mma")?;
66            let mma = ();
67            stream.expect_complete()?;
68            stream.expect_string(".sync")?;
69            let sync = ();
70            stream.expect_complete()?;
71            stream.expect_string(".aligned")?;
72            let aligned = ();
73            stream.expect_complete()?;
74            stream.expect_string(".alayout")?;
75            let alayout = ();
76            stream.expect_complete()?;
77            stream.expect_string(".blayout")?;
78            let blayout = ();
79            stream.expect_complete()?;
80            stream.expect_string(".shape")?;
81            let shape = ();
82            stream.expect_complete()?;
83            stream.expect_string(".dtype")?;
84            let dtype = ();
85            stream.expect_complete()?;
86            stream.expect_string(".ctype")?;
87            let ctype = ();
88            stream.expect_complete()?;
89            let d = GeneralOperand::parse(stream)?;
90            stream.expect_complete()?;
91            stream.expect(&PtxToken::Comma)?;
92            let a = GeneralOperand::parse(stream)?;
93            stream.expect_complete()?;
94            stream.expect(&PtxToken::Comma)?;
95            let b = GeneralOperand::parse(stream)?;
96            stream.expect_complete()?;
97            stream.expect(&PtxToken::Comma)?;
98            let c = GeneralOperand::parse(stream)?;
99            stream.expect_complete()?;
100            stream.expect_complete()?;
101            stream.expect(&PtxToken::Semicolon)?;
102            Ok(WmmaMmaSyncAlignedAlayoutBlayoutShapeDtypeCtype {
103                mma,
104                sync,
105                aligned,
106                alayout,
107                blayout,
108                shape,
109                dtype,
110                ctype,
111                d,
112                a,
113                b,
114                c,
115            })
116        }
117    }
118
119
120}
121
122pub mod section_1 {
123    use super::*;
124    use crate::r#type::instruction::wmma_mma::section_1::*;
125
126    // ============================================================================
127    // Generated enum parsers
128    // ============================================================================
129
130    impl PtxParser for Alayout {
131        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
132            // Try Row
133            {
134                let saved_pos = stream.position();
135                if stream.expect_string(".row").is_ok() {
136                    return Ok(Alayout::Row);
137                }
138                stream.set_position(saved_pos);
139            }
140            let saved_pos = stream.position();
141            // Try Col
142            {
143                let saved_pos = stream.position();
144                if stream.expect_string(".col").is_ok() {
145                    return Ok(Alayout::Col);
146                }
147                stream.set_position(saved_pos);
148            }
149            stream.set_position(saved_pos);
150            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
151            let expected = &[".row", ".col"];
152            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
153            Err(crate::parser::unexpected_value(span, expected, found))
154        }
155    }
156
157    impl PtxParser for Atype {
158        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
159            // Try S8
160            {
161                let saved_pos = stream.position();
162                if stream.expect_string(".s8").is_ok() {
163                    return Ok(Atype::S8);
164                }
165                stream.set_position(saved_pos);
166            }
167            let saved_pos = stream.position();
168            // Try U8
169            {
170                let saved_pos = stream.position();
171                if stream.expect_string(".u8").is_ok() {
172                    return Ok(Atype::U8);
173                }
174                stream.set_position(saved_pos);
175            }
176            stream.set_position(saved_pos);
177            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
178            let expected = &[".s8", ".u8"];
179            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
180            Err(crate::parser::unexpected_value(span, expected, found))
181        }
182    }
183
184    impl PtxParser for Blayout {
185        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
186            // Try Row
187            {
188                let saved_pos = stream.position();
189                if stream.expect_string(".row").is_ok() {
190                    return Ok(Blayout::Row);
191                }
192                stream.set_position(saved_pos);
193            }
194            let saved_pos = stream.position();
195            // Try Col
196            {
197                let saved_pos = stream.position();
198                if stream.expect_string(".col").is_ok() {
199                    return Ok(Blayout::Col);
200                }
201                stream.set_position(saved_pos);
202            }
203            stream.set_position(saved_pos);
204            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
205            let expected = &[".row", ".col"];
206            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
207            Err(crate::parser::unexpected_value(span, expected, found))
208        }
209    }
210
211    impl PtxParser for Btype {
212        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
213            // Try S8
214            {
215                let saved_pos = stream.position();
216                if stream.expect_string(".s8").is_ok() {
217                    return Ok(Btype::S8);
218                }
219                stream.set_position(saved_pos);
220            }
221            let saved_pos = stream.position();
222            // Try U8
223            {
224                let saved_pos = stream.position();
225                if stream.expect_string(".u8").is_ok() {
226                    return Ok(Btype::U8);
227                }
228                stream.set_position(saved_pos);
229            }
230            stream.set_position(saved_pos);
231            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
232            let expected = &[".s8", ".u8"];
233            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
234            Err(crate::parser::unexpected_value(span, expected, found))
235        }
236    }
237
238    impl PtxParser for Shape {
239        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
240            // Try M16n16k16
241            {
242                let saved_pos = stream.position();
243                if stream.expect_string(".m16n16k16").is_ok() {
244                    return Ok(Shape::M16n16k16);
245                }
246                stream.set_position(saved_pos);
247            }
248            let saved_pos = stream.position();
249            // Try M8n32k16
250            {
251                let saved_pos = stream.position();
252                if stream.expect_string(".m8n32k16").is_ok() {
253                    return Ok(Shape::M8n32k16);
254                }
255                stream.set_position(saved_pos);
256            }
257            stream.set_position(saved_pos);
258            let saved_pos = stream.position();
259            // Try M32n8k16
260            {
261                let saved_pos = stream.position();
262                if stream.expect_string(".m32n8k16").is_ok() {
263                    return Ok(Shape::M32n8k16);
264                }
265                stream.set_position(saved_pos);
266            }
267            stream.set_position(saved_pos);
268            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
269            let expected = &[".m16n16k16", ".m8n32k16", ".m32n8k16"];
270            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
271            Err(crate::parser::unexpected_value(span, expected, found))
272        }
273    }
274
275    impl PtxParser for WmmaMmaSyncAlignedAlayoutBlayoutShapeS32AtypeBtypeS32Satfinite {
276        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
277            stream.expect_string("wmma")?;
278            stream.expect_string(".mma")?;
279            let mma = ();
280            stream.expect_complete()?;
281            stream.expect_string(".sync")?;
282            let sync = ();
283            stream.expect_complete()?;
284            stream.expect_string(".aligned")?;
285            let aligned = ();
286            stream.expect_complete()?;
287            let alayout = Alayout::parse(stream)?;
288            stream.expect_complete()?;
289            let blayout = Blayout::parse(stream)?;
290            stream.expect_complete()?;
291            let shape = Shape::parse(stream)?;
292            stream.expect_complete()?;
293            stream.expect_string(".s32")?;
294            let s32 = ();
295            stream.expect_complete()?;
296            let atype = Atype::parse(stream)?;
297            stream.expect_complete()?;
298            let btype = Btype::parse(stream)?;
299            stream.expect_complete()?;
300            stream.expect_string(".s32")?;
301            let s322 = ();
302            stream.expect_complete()?;
303            let saved_pos = stream.position();
304            let satfinite = stream.expect_string(".satfinite").is_ok();
305            if !satfinite {
306                stream.set_position(saved_pos);
307            }
308            stream.expect_complete()?;
309            let d = GeneralOperand::parse(stream)?;
310            stream.expect_complete()?;
311            stream.expect(&PtxToken::Comma)?;
312            let a = GeneralOperand::parse(stream)?;
313            stream.expect_complete()?;
314            stream.expect(&PtxToken::Comma)?;
315            let b = GeneralOperand::parse(stream)?;
316            stream.expect_complete()?;
317            stream.expect(&PtxToken::Comma)?;
318            let c = GeneralOperand::parse(stream)?;
319            stream.expect_complete()?;
320            stream.expect_complete()?;
321            stream.expect(&PtxToken::Semicolon)?;
322            Ok(WmmaMmaSyncAlignedAlayoutBlayoutShapeS32AtypeBtypeS32Satfinite {
323                mma,
324                sync,
325                aligned,
326                alayout,
327                blayout,
328                shape,
329                s32,
330                atype,
331                btype,
332                s322,
333                satfinite,
334                d,
335                a,
336                b,
337                c,
338            })
339        }
340    }
341
342
343}
344
345pub mod section_2 {
346    use super::*;
347    use crate::r#type::instruction::wmma_mma::section_2::*;
348
349    // ============================================================================
350    // Generated enum parsers
351    // ============================================================================
352
353    impl PtxParser for Alayout {
354        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
355            // Try Row
356            {
357                let saved_pos = stream.position();
358                if stream.expect_string(".row").is_ok() {
359                    return Ok(Alayout::Row);
360                }
361                stream.set_position(saved_pos);
362            }
363            let saved_pos = stream.position();
364            // Try Col
365            {
366                let saved_pos = stream.position();
367                if stream.expect_string(".col").is_ok() {
368                    return Ok(Alayout::Col);
369                }
370                stream.set_position(saved_pos);
371            }
372            stream.set_position(saved_pos);
373            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
374            let expected = &[".row", ".col"];
375            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
376            Err(crate::parser::unexpected_value(span, expected, found))
377        }
378    }
379
380    impl PtxParser for Atype {
381        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
382            // Try Bf16
383            {
384                let saved_pos = stream.position();
385                if stream.expect_string(".bf16").is_ok() {
386                    return Ok(Atype::Bf16);
387                }
388                stream.set_position(saved_pos);
389            }
390            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
391            let expected = &[".bf16"];
392            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
393            Err(crate::parser::unexpected_value(span, expected, found))
394        }
395    }
396
397    impl PtxParser for Blayout {
398        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
399            // Try Row
400            {
401                let saved_pos = stream.position();
402                if stream.expect_string(".row").is_ok() {
403                    return Ok(Blayout::Row);
404                }
405                stream.set_position(saved_pos);
406            }
407            let saved_pos = stream.position();
408            // Try Col
409            {
410                let saved_pos = stream.position();
411                if stream.expect_string(".col").is_ok() {
412                    return Ok(Blayout::Col);
413                }
414                stream.set_position(saved_pos);
415            }
416            stream.set_position(saved_pos);
417            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
418            let expected = &[".row", ".col"];
419            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
420            Err(crate::parser::unexpected_value(span, expected, found))
421        }
422    }
423
424    impl PtxParser for Btype {
425        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
426            // Try Bf16
427            {
428                let saved_pos = stream.position();
429                if stream.expect_string(".bf16").is_ok() {
430                    return Ok(Btype::Bf16);
431                }
432                stream.set_position(saved_pos);
433            }
434            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
435            let expected = &[".bf16"];
436            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
437            Err(crate::parser::unexpected_value(span, expected, found))
438        }
439    }
440
441    impl PtxParser for Shape {
442        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
443            // Try M16n16k16
444            {
445                let saved_pos = stream.position();
446                if stream.expect_string(".m16n16k16").is_ok() {
447                    return Ok(Shape::M16n16k16);
448                }
449                stream.set_position(saved_pos);
450            }
451            let saved_pos = stream.position();
452            // Try M8n32k16
453            {
454                let saved_pos = stream.position();
455                if stream.expect_string(".m8n32k16").is_ok() {
456                    return Ok(Shape::M8n32k16);
457                }
458                stream.set_position(saved_pos);
459            }
460            stream.set_position(saved_pos);
461            let saved_pos = stream.position();
462            // Try M32n8k16
463            {
464                let saved_pos = stream.position();
465                if stream.expect_string(".m32n8k16").is_ok() {
466                    return Ok(Shape::M32n8k16);
467                }
468                stream.set_position(saved_pos);
469            }
470            stream.set_position(saved_pos);
471            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
472            let expected = &[".m16n16k16", ".m8n32k16", ".m32n8k16"];
473            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
474            Err(crate::parser::unexpected_value(span, expected, found))
475        }
476    }
477
478    impl PtxParser for WmmaMmaSyncAlignedAlayoutBlayoutShapeF32AtypeBtypeF32 {
479        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
480            stream.expect_string("wmma")?;
481            stream.expect_string(".mma")?;
482            let mma = ();
483            stream.expect_complete()?;
484            stream.expect_string(".sync")?;
485            let sync = ();
486            stream.expect_complete()?;
487            stream.expect_string(".aligned")?;
488            let aligned = ();
489            stream.expect_complete()?;
490            let alayout = Alayout::parse(stream)?;
491            stream.expect_complete()?;
492            let blayout = Blayout::parse(stream)?;
493            stream.expect_complete()?;
494            let shape = Shape::parse(stream)?;
495            stream.expect_complete()?;
496            stream.expect_string(".f32")?;
497            let f32 = ();
498            stream.expect_complete()?;
499            let atype = Atype::parse(stream)?;
500            stream.expect_complete()?;
501            let btype = Btype::parse(stream)?;
502            stream.expect_complete()?;
503            stream.expect_string(".f32")?;
504            let f322 = ();
505            stream.expect_complete()?;
506            let d = GeneralOperand::parse(stream)?;
507            stream.expect_complete()?;
508            stream.expect(&PtxToken::Comma)?;
509            let a = GeneralOperand::parse(stream)?;
510            stream.expect_complete()?;
511            stream.expect(&PtxToken::Comma)?;
512            let b = GeneralOperand::parse(stream)?;
513            stream.expect_complete()?;
514            stream.expect(&PtxToken::Comma)?;
515            let c = GeneralOperand::parse(stream)?;
516            stream.expect_complete()?;
517            stream.expect_complete()?;
518            stream.expect(&PtxToken::Semicolon)?;
519            Ok(WmmaMmaSyncAlignedAlayoutBlayoutShapeF32AtypeBtypeF32 {
520                mma,
521                sync,
522                aligned,
523                alayout,
524                blayout,
525                shape,
526                f32,
527                atype,
528                btype,
529                f322,
530                d,
531                a,
532                b,
533                c,
534            })
535        }
536    }
537
538
539}
540
541pub mod section_3 {
542    use super::*;
543    use crate::r#type::instruction::wmma_mma::section_3::*;
544
545    // ============================================================================
546    // Generated enum parsers
547    // ============================================================================
548
549    impl PtxParser for Alayout {
550        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
551            // Try Row
552            {
553                let saved_pos = stream.position();
554                if stream.expect_string(".row").is_ok() {
555                    return Ok(Alayout::Row);
556                }
557                stream.set_position(saved_pos);
558            }
559            let saved_pos = stream.position();
560            // Try Col
561            {
562                let saved_pos = stream.position();
563                if stream.expect_string(".col").is_ok() {
564                    return Ok(Alayout::Col);
565                }
566                stream.set_position(saved_pos);
567            }
568            stream.set_position(saved_pos);
569            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
570            let expected = &[".row", ".col"];
571            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
572            Err(crate::parser::unexpected_value(span, expected, found))
573        }
574    }
575
576    impl PtxParser for Atype {
577        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
578            // Try Tf32
579            {
580                let saved_pos = stream.position();
581                if stream.expect_string(".tf32").is_ok() {
582                    return Ok(Atype::Tf32);
583                }
584                stream.set_position(saved_pos);
585            }
586            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
587            let expected = &[".tf32"];
588            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
589            Err(crate::parser::unexpected_value(span, expected, found))
590        }
591    }
592
593    impl PtxParser for Blayout {
594        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
595            // Try Row
596            {
597                let saved_pos = stream.position();
598                if stream.expect_string(".row").is_ok() {
599                    return Ok(Blayout::Row);
600                }
601                stream.set_position(saved_pos);
602            }
603            let saved_pos = stream.position();
604            // Try Col
605            {
606                let saved_pos = stream.position();
607                if stream.expect_string(".col").is_ok() {
608                    return Ok(Blayout::Col);
609                }
610                stream.set_position(saved_pos);
611            }
612            stream.set_position(saved_pos);
613            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
614            let expected = &[".row", ".col"];
615            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
616            Err(crate::parser::unexpected_value(span, expected, found))
617        }
618    }
619
620    impl PtxParser for Btype {
621        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
622            // Try Tf32
623            {
624                let saved_pos = stream.position();
625                if stream.expect_string(".tf32").is_ok() {
626                    return Ok(Btype::Tf32);
627                }
628                stream.set_position(saved_pos);
629            }
630            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
631            let expected = &[".tf32"];
632            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
633            Err(crate::parser::unexpected_value(span, expected, found))
634        }
635    }
636
637    impl PtxParser for Shape {
638        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
639            // Try M16n16k8
640            {
641                let saved_pos = stream.position();
642                if stream.expect_string(".m16n16k8").is_ok() {
643                    return Ok(Shape::M16n16k8);
644                }
645                stream.set_position(saved_pos);
646            }
647            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
648            let expected = &[".m16n16k8"];
649            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
650            Err(crate::parser::unexpected_value(span, expected, found))
651        }
652    }
653
654    impl PtxParser for WmmaMmaSyncAlignedAlayoutBlayoutShapeF32AtypeBtypeF321 {
655        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
656            stream.expect_string("wmma")?;
657            stream.expect_string(".mma")?;
658            let mma = ();
659            stream.expect_complete()?;
660            stream.expect_string(".sync")?;
661            let sync = ();
662            stream.expect_complete()?;
663            stream.expect_string(".aligned")?;
664            let aligned = ();
665            stream.expect_complete()?;
666            let alayout = Alayout::parse(stream)?;
667            stream.expect_complete()?;
668            let blayout = Blayout::parse(stream)?;
669            stream.expect_complete()?;
670            let shape = Shape::parse(stream)?;
671            stream.expect_complete()?;
672            stream.expect_string(".f32")?;
673            let f32 = ();
674            stream.expect_complete()?;
675            let atype = Atype::parse(stream)?;
676            stream.expect_complete()?;
677            let btype = Btype::parse(stream)?;
678            stream.expect_complete()?;
679            stream.expect_string(".f32")?;
680            let f322 = ();
681            stream.expect_complete()?;
682            let d = GeneralOperand::parse(stream)?;
683            stream.expect_complete()?;
684            stream.expect(&PtxToken::Comma)?;
685            let a = GeneralOperand::parse(stream)?;
686            stream.expect_complete()?;
687            stream.expect(&PtxToken::Comma)?;
688            let b = GeneralOperand::parse(stream)?;
689            stream.expect_complete()?;
690            stream.expect(&PtxToken::Comma)?;
691            let c = GeneralOperand::parse(stream)?;
692            stream.expect_complete()?;
693            stream.expect_complete()?;
694            stream.expect(&PtxToken::Semicolon)?;
695            Ok(WmmaMmaSyncAlignedAlayoutBlayoutShapeF32AtypeBtypeF321 {
696                mma,
697                sync,
698                aligned,
699                alayout,
700                blayout,
701                shape,
702                f32,
703                atype,
704                btype,
705                f322,
706                d,
707                a,
708                b,
709                c,
710            })
711        }
712    }
713
714
715}
716
717pub mod section_4 {
718    use super::*;
719    use crate::r#type::instruction::wmma_mma::section_4::*;
720
721    // ============================================================================
722    // Generated enum parsers
723    // ============================================================================
724
725    impl PtxParser for Alayout {
726        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
727            // Try Row
728            {
729                let saved_pos = stream.position();
730                if stream.expect_string(".row").is_ok() {
731                    return Ok(Alayout::Row);
732                }
733                stream.set_position(saved_pos);
734            }
735            let saved_pos = stream.position();
736            // Try Col
737            {
738                let saved_pos = stream.position();
739                if stream.expect_string(".col").is_ok() {
740                    return Ok(Alayout::Col);
741                }
742                stream.set_position(saved_pos);
743            }
744            stream.set_position(saved_pos);
745            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
746            let expected = &[".row", ".col"];
747            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
748            Err(crate::parser::unexpected_value(span, expected, found))
749        }
750    }
751
752    impl PtxParser for Blayout {
753        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
754            // Try Row
755            {
756                let saved_pos = stream.position();
757                if stream.expect_string(".row").is_ok() {
758                    return Ok(Blayout::Row);
759                }
760                stream.set_position(saved_pos);
761            }
762            let saved_pos = stream.position();
763            // Try Col
764            {
765                let saved_pos = stream.position();
766                if stream.expect_string(".col").is_ok() {
767                    return Ok(Blayout::Col);
768                }
769                stream.set_position(saved_pos);
770            }
771            stream.set_position(saved_pos);
772            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
773            let expected = &[".row", ".col"];
774            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
775            Err(crate::parser::unexpected_value(span, expected, found))
776        }
777    }
778
779    impl PtxParser for Rnd {
780        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
781            // Try Rn
782            {
783                let saved_pos = stream.position();
784                if stream.expect_string(".rn").is_ok() {
785                    return Ok(Rnd::Rn);
786                }
787                stream.set_position(saved_pos);
788            }
789            let saved_pos = stream.position();
790            // Try Rz
791            {
792                let saved_pos = stream.position();
793                if stream.expect_string(".rz").is_ok() {
794                    return Ok(Rnd::Rz);
795                }
796                stream.set_position(saved_pos);
797            }
798            stream.set_position(saved_pos);
799            let saved_pos = stream.position();
800            // Try Rm
801            {
802                let saved_pos = stream.position();
803                if stream.expect_string(".rm").is_ok() {
804                    return Ok(Rnd::Rm);
805                }
806                stream.set_position(saved_pos);
807            }
808            stream.set_position(saved_pos);
809            let saved_pos = stream.position();
810            // Try Rp
811            {
812                let saved_pos = stream.position();
813                if stream.expect_string(".rp").is_ok() {
814                    return Ok(Rnd::Rp);
815                }
816                stream.set_position(saved_pos);
817            }
818            stream.set_position(saved_pos);
819            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
820            let expected = &[".rn", ".rz", ".rm", ".rp"];
821            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
822            Err(crate::parser::unexpected_value(span, expected, found))
823        }
824    }
825
826    impl PtxParser for Shape {
827        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
828            // Try M8n8k4
829            {
830                let saved_pos = stream.position();
831                if stream.expect_string(".m8n8k4").is_ok() {
832                    return Ok(Shape::M8n8k4);
833                }
834                stream.set_position(saved_pos);
835            }
836            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
837            let expected = &[".m8n8k4"];
838            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
839            Err(crate::parser::unexpected_value(span, expected, found))
840        }
841    }
842
843    impl PtxParser for WmmaMmaSyncAlignedAlayoutBlayoutShapeRndF64F64F64F64 {
844        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
845            stream.expect_string("wmma")?;
846            stream.expect_string(".mma")?;
847            let mma = ();
848            stream.expect_complete()?;
849            stream.expect_string(".sync")?;
850            let sync = ();
851            stream.expect_complete()?;
852            stream.expect_string(".aligned")?;
853            let aligned = ();
854            stream.expect_complete()?;
855            let alayout = Alayout::parse(stream)?;
856            stream.expect_complete()?;
857            let blayout = Blayout::parse(stream)?;
858            stream.expect_complete()?;
859            let shape = Shape::parse(stream)?;
860            stream.expect_complete()?;
861            let saved_pos = stream.position();
862            let rnd = match Rnd::parse(stream) {
863                Ok(val) => Some(val),
864                Err(_) => {
865                    stream.set_position(saved_pos);
866                    None
867                }
868            };
869            stream.expect_complete()?;
870            stream.expect_string(".f64")?;
871            let f64 = ();
872            stream.expect_complete()?;
873            stream.expect_string(".f64")?;
874            let f642 = ();
875            stream.expect_complete()?;
876            stream.expect_string(".f64")?;
877            let f644 = ();
878            stream.expect_complete()?;
879            stream.expect_string(".f64")?;
880            let f646 = ();
881            stream.expect_complete()?;
882            let d = GeneralOperand::parse(stream)?;
883            stream.expect_complete()?;
884            stream.expect(&PtxToken::Comma)?;
885            let a = GeneralOperand::parse(stream)?;
886            stream.expect_complete()?;
887            stream.expect(&PtxToken::Comma)?;
888            let b = GeneralOperand::parse(stream)?;
889            stream.expect_complete()?;
890            stream.expect(&PtxToken::Comma)?;
891            let c = GeneralOperand::parse(stream)?;
892            stream.expect_complete()?;
893            stream.expect_complete()?;
894            stream.expect(&PtxToken::Semicolon)?;
895            Ok(WmmaMmaSyncAlignedAlayoutBlayoutShapeRndF64F64F64F64 {
896                mma,
897                sync,
898                aligned,
899                alayout,
900                blayout,
901                shape,
902                rnd,
903                f64,
904                f642,
905                f644,
906                f646,
907                d,
908                a,
909                b,
910                c,
911            })
912        }
913    }
914
915
916}
917
918pub mod section_5 {
919    use super::*;
920    use crate::r#type::instruction::wmma_mma::section_5::*;
921
922    // ============================================================================
923    // Generated enum parsers
924    // ============================================================================
925
926    impl PtxParser for Atype {
927        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
928            // Try S4
929            {
930                let saved_pos = stream.position();
931                if stream.expect_string(".s4").is_ok() {
932                    return Ok(Atype::S4);
933                }
934                stream.set_position(saved_pos);
935            }
936            let saved_pos = stream.position();
937            // Try U4
938            {
939                let saved_pos = stream.position();
940                if stream.expect_string(".u4").is_ok() {
941                    return Ok(Atype::U4);
942                }
943                stream.set_position(saved_pos);
944            }
945            stream.set_position(saved_pos);
946            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
947            let expected = &[".s4", ".u4"];
948            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
949            Err(crate::parser::unexpected_value(span, expected, found))
950        }
951    }
952
953    impl PtxParser for Btype {
954        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
955            // Try S4
956            {
957                let saved_pos = stream.position();
958                if stream.expect_string(".s4").is_ok() {
959                    return Ok(Btype::S4);
960                }
961                stream.set_position(saved_pos);
962            }
963            let saved_pos = stream.position();
964            // Try U4
965            {
966                let saved_pos = stream.position();
967                if stream.expect_string(".u4").is_ok() {
968                    return Ok(Btype::U4);
969                }
970                stream.set_position(saved_pos);
971            }
972            stream.set_position(saved_pos);
973            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
974            let expected = &[".s4", ".u4"];
975            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
976            Err(crate::parser::unexpected_value(span, expected, found))
977        }
978    }
979
980    impl PtxParser for Shape {
981        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
982            // Try M8n8k32
983            {
984                let saved_pos = stream.position();
985                if stream.expect_string(".m8n8k32").is_ok() {
986                    return Ok(Shape::M8n8k32);
987                }
988                stream.set_position(saved_pos);
989            }
990            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
991            let expected = &[".m8n8k32"];
992            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
993            Err(crate::parser::unexpected_value(span, expected, found))
994        }
995    }
996
997    impl PtxParser for WmmaMmaSyncAlignedRowColShapeS32AtypeBtypeS32Satfinite {
998        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
999            stream.expect_string("wmma")?;
1000            stream.expect_string(".mma")?;
1001            let mma = ();
1002            stream.expect_complete()?;
1003            stream.expect_string(".sync")?;
1004            let sync = ();
1005            stream.expect_complete()?;
1006            stream.expect_string(".aligned")?;
1007            let aligned = ();
1008            stream.expect_complete()?;
1009            stream.expect_string(".row")?;
1010            let row = ();
1011            stream.expect_complete()?;
1012            stream.expect_string(".col")?;
1013            let col = ();
1014            stream.expect_complete()?;
1015            let shape = Shape::parse(stream)?;
1016            stream.expect_complete()?;
1017            stream.expect_string(".s32")?;
1018            let s32 = ();
1019            stream.expect_complete()?;
1020            let atype = Atype::parse(stream)?;
1021            stream.expect_complete()?;
1022            let btype = Btype::parse(stream)?;
1023            stream.expect_complete()?;
1024            stream.expect_string(".s32")?;
1025            let s322 = ();
1026            stream.expect_complete()?;
1027            let saved_pos = stream.position();
1028            let satfinite = stream.expect_string(".satfinite").is_ok();
1029            if !satfinite {
1030                stream.set_position(saved_pos);
1031            }
1032            stream.expect_complete()?;
1033            let d = GeneralOperand::parse(stream)?;
1034            stream.expect_complete()?;
1035            stream.expect(&PtxToken::Comma)?;
1036            let a = GeneralOperand::parse(stream)?;
1037            stream.expect_complete()?;
1038            stream.expect(&PtxToken::Comma)?;
1039            let b = GeneralOperand::parse(stream)?;
1040            stream.expect_complete()?;
1041            stream.expect(&PtxToken::Comma)?;
1042            let c = GeneralOperand::parse(stream)?;
1043            stream.expect_complete()?;
1044            stream.expect_complete()?;
1045            stream.expect(&PtxToken::Semicolon)?;
1046            Ok(WmmaMmaSyncAlignedRowColShapeS32AtypeBtypeS32Satfinite {
1047                mma,
1048                sync,
1049                aligned,
1050                row,
1051                col,
1052                shape,
1053                s32,
1054                atype,
1055                btype,
1056                s322,
1057                satfinite,
1058                d,
1059                a,
1060                b,
1061                c,
1062            })
1063        }
1064    }
1065
1066
1067}
1068
1069pub mod section_6 {
1070    use super::*;
1071    use crate::r#type::instruction::wmma_mma::section_6::*;
1072
1073    // ============================================================================
1074    // Generated enum parsers
1075    // ============================================================================
1076
1077    impl PtxParser for Atype {
1078        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1079            // Try B1
1080            {
1081                let saved_pos = stream.position();
1082                if stream.expect_string(".b1").is_ok() {
1083                    return Ok(Atype::B1);
1084                }
1085                stream.set_position(saved_pos);
1086            }
1087            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1088            let expected = &[".b1"];
1089            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1090            Err(crate::parser::unexpected_value(span, expected, found))
1091        }
1092    }
1093
1094    impl PtxParser for Btype {
1095        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1096            // Try B1
1097            {
1098                let saved_pos = stream.position();
1099                if stream.expect_string(".b1").is_ok() {
1100                    return Ok(Btype::B1);
1101                }
1102                stream.set_position(saved_pos);
1103            }
1104            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1105            let expected = &[".b1"];
1106            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1107            Err(crate::parser::unexpected_value(span, expected, found))
1108        }
1109    }
1110
1111    impl PtxParser for Op {
1112        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1113            // Try Xor
1114            {
1115                let saved_pos = stream.position();
1116                if stream.expect_string(".xor").is_ok() {
1117                    return Ok(Op::Xor);
1118                }
1119                stream.set_position(saved_pos);
1120            }
1121            let saved_pos = stream.position();
1122            // Try And
1123            {
1124                let saved_pos = stream.position();
1125                if stream.expect_string(".and").is_ok() {
1126                    return Ok(Op::And);
1127                }
1128                stream.set_position(saved_pos);
1129            }
1130            stream.set_position(saved_pos);
1131            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1132            let expected = &[".xor", ".and"];
1133            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1134            Err(crate::parser::unexpected_value(span, expected, found))
1135        }
1136    }
1137
1138    impl PtxParser for Shape {
1139        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1140            // Try M8n8k128
1141            {
1142                let saved_pos = stream.position();
1143                if stream.expect_string(".m8n8k128").is_ok() {
1144                    return Ok(Shape::M8n8k128);
1145                }
1146                stream.set_position(saved_pos);
1147            }
1148            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1149            let expected = &[".m8n8k128"];
1150            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1151            Err(crate::parser::unexpected_value(span, expected, found))
1152        }
1153    }
1154
1155    impl PtxParser for WmmaMmaOpPopcSyncAlignedRowColShapeS32AtypeBtypeS32 {
1156        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1157            stream.expect_string("wmma")?;
1158            stream.expect_string(".mma")?;
1159            let mma = ();
1160            stream.expect_complete()?;
1161            let op = Op::parse(stream)?;
1162            stream.expect_complete()?;
1163            stream.expect_string(".popc")?;
1164            let popc = ();
1165            stream.expect_complete()?;
1166            stream.expect_string(".sync")?;
1167            let sync = ();
1168            stream.expect_complete()?;
1169            stream.expect_string(".aligned")?;
1170            let aligned = ();
1171            stream.expect_complete()?;
1172            stream.expect_string(".row")?;
1173            let row = ();
1174            stream.expect_complete()?;
1175            stream.expect_string(".col")?;
1176            let col = ();
1177            stream.expect_complete()?;
1178            let shape = Shape::parse(stream)?;
1179            stream.expect_complete()?;
1180            stream.expect_string(".s32")?;
1181            let s32 = ();
1182            stream.expect_complete()?;
1183            let atype = Atype::parse(stream)?;
1184            stream.expect_complete()?;
1185            let btype = Btype::parse(stream)?;
1186            stream.expect_complete()?;
1187            stream.expect_string(".s32")?;
1188            let s322 = ();
1189            stream.expect_complete()?;
1190            let d = GeneralOperand::parse(stream)?;
1191            stream.expect_complete()?;
1192            stream.expect(&PtxToken::Comma)?;
1193            let a = GeneralOperand::parse(stream)?;
1194            stream.expect_complete()?;
1195            stream.expect(&PtxToken::Comma)?;
1196            let b = GeneralOperand::parse(stream)?;
1197            stream.expect_complete()?;
1198            stream.expect(&PtxToken::Comma)?;
1199            let c = GeneralOperand::parse(stream)?;
1200            stream.expect_complete()?;
1201            stream.expect_complete()?;
1202            stream.expect(&PtxToken::Semicolon)?;
1203            Ok(WmmaMmaOpPopcSyncAlignedRowColShapeS32AtypeBtypeS32 {
1204                mma,
1205                op,
1206                popc,
1207                sync,
1208                aligned,
1209                row,
1210                col,
1211                shape,
1212                s32,
1213                atype,
1214                btype,
1215                s322,
1216                d,
1217                a,
1218                b,
1219                c,
1220            })
1221        }
1222    }
1223
1224
1225}
1226
ptx_parser/parser/instruction/wmma_mma.rs

ptx_parser/parser/instruction/
wmma_mma.rs
