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::parser::{
55    PtxParseError, PtxParser, PtxTokenStream, Span,
56    util::{
57        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
58        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
59    },
60};
61use crate::r#type::common::*;
62use crate::{alt, ok, seq_n};
63
64pub mod section_0 {
65    use super::*;
66    use crate::r#type::instruction::wmma_mma::section_0::*;
67
68    impl PtxParser for WmmaMmaSyncAlignedAlayoutBlayoutShapeDtypeCtype {
69        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
70            try_map(
71                seq_n!(
72                    string_p("wmma"),
73                    string_p(".mma"),
74                    string_p(".sync"),
75                    string_p(".aligned"),
76                    string_p(".alayout"),
77                    string_p(".blayout"),
78                    string_p(".shape"),
79                    string_p(".dtype"),
80                    string_p(".ctype"),
81                    GeneralOperand::parse(),
82                    comma_p(),
83                    GeneralOperand::parse(),
84                    comma_p(),
85                    GeneralOperand::parse(),
86                    comma_p(),
87                    GeneralOperand::parse(),
88                    semicolon_p()
89                ),
90                |(
91                    _,
92                    mma,
93                    sync,
94                    aligned,
95                    alayout,
96                    blayout,
97                    shape,
98                    dtype,
99                    ctype,
100                    d,
101                    _,
102                    a,
103                    _,
104                    b,
105                    _,
106                    c,
107                    _,
108                ),
109                 span| {
110                    ok!(WmmaMmaSyncAlignedAlayoutBlayoutShapeDtypeCtype {
111                        mma = mma,
112                        sync = sync,
113                        aligned = aligned,
114                        alayout = alayout,
115                        blayout = blayout,
116                        shape = shape,
117                        dtype = dtype,
118                        ctype = ctype,
119                        d = d,
120                        a = a,
121                        b = b,
122                        c = c,
123
124                    })
125                },
126            )
127        }
128    }
129}
130
131pub mod section_1 {
132    use super::*;
133    use crate::r#type::instruction::wmma_mma::section_1::*;
134
135    // ============================================================================
136    // Generated enum parsers
137    // ============================================================================
138
139    impl PtxParser for Alayout {
140        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
141            alt!(
142                map(string_p(".row"), |_, _span| Alayout::Row),
143                map(string_p(".col"), |_, _span| Alayout::Col)
144            )
145        }
146    }
147
148    impl PtxParser for Atype {
149        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
150            alt!(
151                map(string_p(".s8"), |_, _span| Atype::S8),
152                map(string_p(".u8"), |_, _span| Atype::U8)
153            )
154        }
155    }
156
157    impl PtxParser for Blayout {
158        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
159            alt!(
160                map(string_p(".row"), |_, _span| Blayout::Row),
161                map(string_p(".col"), |_, _span| Blayout::Col)
162            )
163        }
164    }
165
166    impl PtxParser for Btype {
167        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
168            alt!(
169                map(string_p(".s8"), |_, _span| Btype::S8),
170                map(string_p(".u8"), |_, _span| Btype::U8)
171            )
172        }
173    }
174
175    impl PtxParser for Shape {
176        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
177            alt!(
178                map(string_p(".m16n16k16"), |_, _span| Shape::M16n16k16),
179                map(string_p(".m8n32k16"), |_, _span| Shape::M8n32k16),
180                map(string_p(".m32n8k16"), |_, _span| Shape::M32n8k16)
181            )
182        }
183    }
184
185    impl PtxParser for WmmaMmaSyncAlignedAlayoutBlayoutShapeS32AtypeBtypeS32Satfinite {
186        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
187            try_map(
188                seq_n!(
189                    string_p("wmma"),
190                    string_p(".mma"),
191                    string_p(".sync"),
192                    string_p(".aligned"),
193                    Alayout::parse(),
194                    Blayout::parse(),
195                    Shape::parse(),
196                    string_p(".s32"),
197                    Atype::parse(),
198                    Btype::parse(),
199                    string_p(".s32"),
200                    map(optional(string_p(".satfinite")), |value, _| value.is_some()),
201                    GeneralOperand::parse(),
202                    comma_p(),
203                    GeneralOperand::parse(),
204                    comma_p(),
205                    GeneralOperand::parse(),
206                    comma_p(),
207                    GeneralOperand::parse(),
208                    semicolon_p()
209                ),
210                |(
211                    _,
212                    mma,
213                    sync,
214                    aligned,
215                    alayout,
216                    blayout,
217                    shape,
218                    s32,
219                    atype,
220                    btype,
221                    s322,
222                    satfinite,
223                    d,
224                    _,
225                    a,
226                    _,
227                    b,
228                    _,
229                    c,
230                    _,
231                ),
232                 span| {
233                    ok!(WmmaMmaSyncAlignedAlayoutBlayoutShapeS32AtypeBtypeS32Satfinite {
234                        mma = mma,
235                        sync = sync,
236                        aligned = aligned,
237                        alayout = alayout,
238                        blayout = blayout,
239                        shape = shape,
240                        s32 = s32,
241                        atype = atype,
242                        btype = btype,
243                        s322 = s322,
244                        satfinite = satfinite,
245                        d = d,
246                        a = a,
247                        b = b,
248                        c = c,
249
250                    })
251                },
252            )
253        }
254    }
255}
256
257pub mod section_2 {
258    use super::*;
259    use crate::r#type::instruction::wmma_mma::section_2::*;
260
261    // ============================================================================
262    // Generated enum parsers
263    // ============================================================================
264
265    impl PtxParser for Alayout {
266        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
267            alt!(
268                map(string_p(".row"), |_, _span| Alayout::Row),
269                map(string_p(".col"), |_, _span| Alayout::Col)
270            )
271        }
272    }
273
274    impl PtxParser for Atype {
275        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
276            alt!(map(string_p(".bf16"), |_, _span| Atype::Bf16))
277        }
278    }
279
280    impl PtxParser for Blayout {
281        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
282            alt!(
283                map(string_p(".row"), |_, _span| Blayout::Row),
284                map(string_p(".col"), |_, _span| Blayout::Col)
285            )
286        }
287    }
288
289    impl PtxParser for Btype {
290        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
291            alt!(map(string_p(".bf16"), |_, _span| Btype::Bf16))
292        }
293    }
294
295    impl PtxParser for Shape {
296        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
297            alt!(
298                map(string_p(".m16n16k16"), |_, _span| Shape::M16n16k16),
299                map(string_p(".m8n32k16"), |_, _span| Shape::M8n32k16),
300                map(string_p(".m32n8k16"), |_, _span| Shape::M32n8k16)
301            )
302        }
303    }
304
305    impl PtxParser for WmmaMmaSyncAlignedAlayoutBlayoutShapeF32AtypeBtypeF32 {
306        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
307            try_map(
308                seq_n!(
309                    string_p("wmma"),
310                    string_p(".mma"),
311                    string_p(".sync"),
312                    string_p(".aligned"),
313                    Alayout::parse(),
314                    Blayout::parse(),
315                    Shape::parse(),
316                    string_p(".f32"),
317                    Atype::parse(),
318                    Btype::parse(),
319                    string_p(".f32"),
320                    GeneralOperand::parse(),
321                    comma_p(),
322                    GeneralOperand::parse(),
323                    comma_p(),
324                    GeneralOperand::parse(),
325                    comma_p(),
326                    GeneralOperand::parse(),
327                    semicolon_p()
328                ),
329                |(
330                    _,
331                    mma,
332                    sync,
333                    aligned,
334                    alayout,
335                    blayout,
336                    shape,
337                    f32,
338                    atype,
339                    btype,
340                    f322,
341                    d,
342                    _,
343                    a,
344                    _,
345                    b,
346                    _,
347                    c,
348                    _,
349                ),
350                 span| {
351                    ok!(WmmaMmaSyncAlignedAlayoutBlayoutShapeF32AtypeBtypeF32 {
352                        mma = mma,
353                        sync = sync,
354                        aligned = aligned,
355                        alayout = alayout,
356                        blayout = blayout,
357                        shape = shape,
358                        f32 = f32,
359                        atype = atype,
360                        btype = btype,
361                        f322 = f322,
362                        d = d,
363                        a = a,
364                        b = b,
365                        c = c,
366
367                    })
368                },
369            )
370        }
371    }
372}
373
374pub mod section_3 {
375    use super::*;
376    use crate::r#type::instruction::wmma_mma::section_3::*;
377
378    // ============================================================================
379    // Generated enum parsers
380    // ============================================================================
381
382    impl PtxParser for Alayout {
383        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
384            alt!(
385                map(string_p(".row"), |_, _span| Alayout::Row),
386                map(string_p(".col"), |_, _span| Alayout::Col)
387            )
388        }
389    }
390
391    impl PtxParser for Atype {
392        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
393            alt!(map(string_p(".tf32"), |_, _span| Atype::Tf32))
394        }
395    }
396
397    impl PtxParser for Blayout {
398        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
399            alt!(
400                map(string_p(".row"), |_, _span| Blayout::Row),
401                map(string_p(".col"), |_, _span| Blayout::Col)
402            )
403        }
404    }
405
406    impl PtxParser for Btype {
407        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
408            alt!(map(string_p(".tf32"), |_, _span| Btype::Tf32))
409        }
410    }
411
412    impl PtxParser for Shape {
413        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
414            alt!(map(string_p(".m16n16k8"), |_, _span| Shape::M16n16k8))
415        }
416    }
417
418    impl PtxParser for WmmaMmaSyncAlignedAlayoutBlayoutShapeF32AtypeBtypeF321 {
419        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
420            try_map(
421                seq_n!(
422                    string_p("wmma"),
423                    string_p(".mma"),
424                    string_p(".sync"),
425                    string_p(".aligned"),
426                    Alayout::parse(),
427                    Blayout::parse(),
428                    Shape::parse(),
429                    string_p(".f32"),
430                    Atype::parse(),
431                    Btype::parse(),
432                    string_p(".f32"),
433                    GeneralOperand::parse(),
434                    comma_p(),
435                    GeneralOperand::parse(),
436                    comma_p(),
437                    GeneralOperand::parse(),
438                    comma_p(),
439                    GeneralOperand::parse(),
440                    semicolon_p()
441                ),
442                |(
443                    _,
444                    mma,
445                    sync,
446                    aligned,
447                    alayout,
448                    blayout,
449                    shape,
450                    f32,
451                    atype,
452                    btype,
453                    f322,
454                    d,
455                    _,
456                    a,
457                    _,
458                    b,
459                    _,
460                    c,
461                    _,
462                ),
463                 span| {
464                    ok!(WmmaMmaSyncAlignedAlayoutBlayoutShapeF32AtypeBtypeF321 {
465                        mma = mma,
466                        sync = sync,
467                        aligned = aligned,
468                        alayout = alayout,
469                        blayout = blayout,
470                        shape = shape,
471                        f32 = f32,
472                        atype = atype,
473                        btype = btype,
474                        f322 = f322,
475                        d = d,
476                        a = a,
477                        b = b,
478                        c = c,
479
480                    })
481                },
482            )
483        }
484    }
485}
486
487pub mod section_4 {
488    use super::*;
489    use crate::r#type::instruction::wmma_mma::section_4::*;
490
491    // ============================================================================
492    // Generated enum parsers
493    // ============================================================================
494
495    impl PtxParser for Alayout {
496        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
497            alt!(
498                map(string_p(".row"), |_, _span| Alayout::Row),
499                map(string_p(".col"), |_, _span| Alayout::Col)
500            )
501        }
502    }
503
504    impl PtxParser for Blayout {
505        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
506            alt!(
507                map(string_p(".row"), |_, _span| Blayout::Row),
508                map(string_p(".col"), |_, _span| Blayout::Col)
509            )
510        }
511    }
512
513    impl PtxParser for Rnd {
514        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
515            alt!(
516                map(string_p(".rn"), |_, _span| Rnd::Rn),
517                map(string_p(".rz"), |_, _span| Rnd::Rz),
518                map(string_p(".rm"), |_, _span| Rnd::Rm),
519                map(string_p(".rp"), |_, _span| Rnd::Rp)
520            )
521        }
522    }
523
524    impl PtxParser for Shape {
525        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
526            alt!(map(string_p(".m8n8k4"), |_, _span| Shape::M8n8k4))
527        }
528    }
529
530    impl PtxParser for WmmaMmaSyncAlignedAlayoutBlayoutShapeRndF64F64F64F64 {
531        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
532            try_map(
533                seq_n!(
534                    string_p("wmma"),
535                    string_p(".mma"),
536                    string_p(".sync"),
537                    string_p(".aligned"),
538                    Alayout::parse(),
539                    Blayout::parse(),
540                    Shape::parse(),
541                    optional(Rnd::parse()),
542                    string_p(".f64"),
543                    string_p(".f64"),
544                    string_p(".f64"),
545                    string_p(".f64"),
546                    GeneralOperand::parse(),
547                    comma_p(),
548                    GeneralOperand::parse(),
549                    comma_p(),
550                    GeneralOperand::parse(),
551                    comma_p(),
552                    GeneralOperand::parse(),
553                    semicolon_p()
554                ),
555                |(
556                    _,
557                    mma,
558                    sync,
559                    aligned,
560                    alayout,
561                    blayout,
562                    shape,
563                    rnd,
564                    f64,
565                    f642,
566                    f644,
567                    f646,
568                    d,
569                    _,
570                    a,
571                    _,
572                    b,
573                    _,
574                    c,
575                    _,
576                ),
577                 span| {
578                    ok!(WmmaMmaSyncAlignedAlayoutBlayoutShapeRndF64F64F64F64 {
579                        mma = mma,
580                        sync = sync,
581                        aligned = aligned,
582                        alayout = alayout,
583                        blayout = blayout,
584                        shape = shape,
585                        rnd = rnd,
586                        f64 = f64,
587                        f642 = f642,
588                        f644 = f644,
589                        f646 = f646,
590                        d = d,
591                        a = a,
592                        b = b,
593                        c = c,
594
595                    })
596                },
597            )
598        }
599    }
600}
601
602pub mod section_5 {
603    use super::*;
604    use crate::r#type::instruction::wmma_mma::section_5::*;
605
606    // ============================================================================
607    // Generated enum parsers
608    // ============================================================================
609
610    impl PtxParser for Atype {
611        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
612            alt!(
613                map(string_p(".s4"), |_, _span| Atype::S4),
614                map(string_p(".u4"), |_, _span| Atype::U4)
615            )
616        }
617    }
618
619    impl PtxParser for Btype {
620        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
621            alt!(
622                map(string_p(".s4"), |_, _span| Btype::S4),
623                map(string_p(".u4"), |_, _span| Btype::U4)
624            )
625        }
626    }
627
628    impl PtxParser for Shape {
629        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
630            alt!(map(string_p(".m8n8k32"), |_, _span| Shape::M8n8k32))
631        }
632    }
633
634    impl PtxParser for WmmaMmaSyncAlignedRowColShapeS32AtypeBtypeS32Satfinite {
635        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
636            try_map(
637                seq_n!(
638                    string_p("wmma"),
639                    string_p(".mma"),
640                    string_p(".sync"),
641                    string_p(".aligned"),
642                    string_p(".row"),
643                    string_p(".col"),
644                    Shape::parse(),
645                    string_p(".s32"),
646                    Atype::parse(),
647                    Btype::parse(),
648                    string_p(".s32"),
649                    map(optional(string_p(".satfinite")), |value, _| value.is_some()),
650                    GeneralOperand::parse(),
651                    comma_p(),
652                    GeneralOperand::parse(),
653                    comma_p(),
654                    GeneralOperand::parse(),
655                    comma_p(),
656                    GeneralOperand::parse(),
657                    semicolon_p()
658                ),
659                |(
660                    _,
661                    mma,
662                    sync,
663                    aligned,
664                    row,
665                    col,
666                    shape,
667                    s32,
668                    atype,
669                    btype,
670                    s322,
671                    satfinite,
672                    d,
673                    _,
674                    a,
675                    _,
676                    b,
677                    _,
678                    c,
679                    _,
680                ),
681                 span| {
682                    ok!(WmmaMmaSyncAlignedRowColShapeS32AtypeBtypeS32Satfinite {
683                        mma = mma,
684                        sync = sync,
685                        aligned = aligned,
686                        row = row,
687                        col = col,
688                        shape = shape,
689                        s32 = s32,
690                        atype = atype,
691                        btype = btype,
692                        s322 = s322,
693                        satfinite = satfinite,
694                        d = d,
695                        a = a,
696                        b = b,
697                        c = c,
698
699                    })
700                },
701            )
702        }
703    }
704}
705
706pub mod section_6 {
707    use super::*;
708    use crate::r#type::instruction::wmma_mma::section_6::*;
709
710    // ============================================================================
711    // Generated enum parsers
712    // ============================================================================
713
714    impl PtxParser for Atype {
715        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
716            alt!(map(string_p(".b1"), |_, _span| Atype::B1))
717        }
718    }
719
720    impl PtxParser for Btype {
721        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
722            alt!(map(string_p(".b1"), |_, _span| Btype::B1))
723        }
724    }
725
726    impl PtxParser for Op {
727        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
728            alt!(
729                map(string_p(".xor"), |_, _span| Op::Xor),
730                map(string_p(".and"), |_, _span| Op::And)
731            )
732        }
733    }
734
735    impl PtxParser for Shape {
736        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
737            alt!(map(string_p(".m8n8k128"), |_, _span| Shape::M8n8k128))
738        }
739    }
740
741    impl PtxParser for WmmaMmaOpPopcSyncAlignedRowColShapeS32AtypeBtypeS32 {
742        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
743            try_map(
744                seq_n!(
745                    string_p("wmma"),
746                    string_p(".mma"),
747                    Op::parse(),
748                    string_p(".popc"),
749                    string_p(".sync"),
750                    string_p(".aligned"),
751                    string_p(".row"),
752                    string_p(".col"),
753                    Shape::parse(),
754                    string_p(".s32"),
755                    Atype::parse(),
756                    Btype::parse(),
757                    string_p(".s32"),
758                    GeneralOperand::parse(),
759                    comma_p(),
760                    GeneralOperand::parse(),
761                    comma_p(),
762                    GeneralOperand::parse(),
763                    comma_p(),
764                    GeneralOperand::parse(),
765                    semicolon_p()
766                ),
767                |(
768                    _,
769                    mma,
770                    op,
771                    popc,
772                    sync,
773                    aligned,
774                    row,
775                    col,
776                    shape,
777                    s32,
778                    atype,
779                    btype,
780                    s322,
781                    d,
782                    _,
783                    a,
784                    _,
785                    b,
786                    _,
787                    c,
788                    _,
789                ),
790                 span| {
791                    ok!(WmmaMmaOpPopcSyncAlignedRowColShapeS32AtypeBtypeS32 {
792                        mma = mma,
793                        op = op,
794                        popc = popc,
795                        sync = sync,
796                        aligned = aligned,
797                        row = row,
798                        col = col,
799                        shape = shape,
800                        s32 = s32,
801                        atype = atype,
802                        btype = btype,
803                        s322 = s322,
804                        d = d,
805                        a = a,
806                        b = b,
807                        c = c,
808
809                    })
810                },
811            )
812        }
813    }
814}
ptx_parser/parser/instruction/wmma_mma.rs

ptx_parser/parser/instruction/
wmma_mma.rs
