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

ptx_parser/parser/instruction/
wmma_mma.rs
