ptx_parser/type/instruction/
mma.rs

1//! Original PTX specification:
2//!
3//! // Half precision floating point type:
4//! mma.sync.aligned.m8n8k4.alayout.blayout.dtype.f16.f16.ctype  d, a, b, c;
5//! mma.sync.aligned.m16n8k8.row.col.dtype.f16.f16.ctype  d, a, b, c;
6//! mma.sync.aligned.m16n8k16.row.col.dtype.f16.f16.ctype d, a, b, c;
7//! .alayout = {.row, .col};
8//! .blayout = {.row, .col};
9//! .ctype   = {.f16, .f32};
10//! .dtype   = {.f16, .f32};
11//! ----------------------------------------------------
12//! // Alternate floating point type:
13//! // Alternate floating point type:
14//! mma.sync.aligned.m16n8k4.row.col.f32.tf32.tf32.f32        d, a, b, c;
15//! mma.sync.aligned.m16n8k8.row.col.f32.atype.btype.f32      d, a, b, c;
16//! mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32       d, a, b, c;
17//! mma.sync.aligned.shape.row.col.dtype.f8type.f8type.ctype  d, a, b, c;
18//! mma.sync.aligned.m16n8k32.row.col.kind.dtype.f8f6f4type.f8f6f4type.ctype d, a, b, c;
19//! .atype      = {.bf16, .tf32};
20//! .btype      = {.bf16, .tf32};
21//! .f8type     = {.e4m3, .e5m2};
22//! .f8f6f4type = {.e4m3, .e5m2, .e3m2, .e2m3, .e2m1};
23//! .ctype      = {.f16, .f32};
24//! .dtype      = {.f16, .f32};
25//! .shape      = {.m16n8k16, .m16n8k32};
26//! .kind       = {.kind::f8f6f4};
27//! ----------------------------------------------------
28//! // Alternate floating point type:
29//! // Alternate floating point type with block scaling:
30//! mma.sync.aligned.m16n8k64.row.col.kind.block_scale{.scale_vec_size}.f32.e2m1.e2m1.f32.stype d, a, b, c, scale-a-data, {byte-id-a, thread-id-a}, scale-b-data, {byte-id-b, thread-id-b};
31//! .kind           = {.kind::mxf4};
32//! .scale_vec_size = {.scale_vec::2X};
33//! .stype          = {.ue8m0};
34//! ----------------------------------------------------
35//! // Alternate floating point type:
36//! mma.sync.aligned.m16n8k64.row.col.kind.block_scale.scale_vec_size.f32.e2m1.e2m1.f32.stype d, a, b, c, scale-a-data, {byte-id-a, thread-id-a}, scale-b-data, {byte-id-b, thread-id-b};
37//! .kind           = {.kind::mxf4nvf4};
38//! .scale_vec_size = {.scale_vec::2X, .scale_vec::4X};
39//! .stype          = {.ue8m0, .ue4m3};
40//! ----------------------------------------------------
41//! // Alternate floating point type:
42//! mma.sync.aligned.m16n8k32.row.col.kind.block_scale{.scale_vec_size}.f32.f8f6f4type.f8f6f4type.f32.stype d, a, b, c, scale-a-data, {byte-id-a, thread-id-a}, scale-b-data, {byte-id-b, thread-id-b};
43//! .kind           = {.kind::mxf8f6f4};
44//! .scale_vec_size = {.scale_vec::1X};
45//! .f8f6f4type     = {.e4m3, .e5m2, .e3m2, .e2m3, .e2m1};
46//! .stype          = {.ue8m0};
47//! ----------------------------------------------------
48//! // Alternate floating point type:
49//! // Double precision floating point type:
50//! mma.sync.aligned.shape.row.col.f64.f64.f64.f64 d, a, b, c;
51//! .shape   = {.m8n84, .m16n8k4, .m16n8k8, .m16n8k16};
52//! ----------------------------------------------------
53//! // Alternate floating point type:
54//! // Integer type:
55//! mma.sync.aligned.shape.row.col{.satfinite}.s32.atype.btype.s32 d, a, b, c;
56//! .shape   = {.m8n8k16, .m16n8k16, .m16n8k32};
57//! .atype   = {.u8, .s8};
58//! .btype   = {.u8, .s8};
59//! ----------------------------------------------------
60//! // Alternate floating point type:
61//! mma.sync.aligned.shape.row.col{.satfinite}.s32.atype.btype.s32 d, a, b, c;
62//! .shape   = {.m8n8k32, .m16n8k32, .m16n8k64};
63//! .atype   = {.u4, .s4};
64//! .btype   = {.u4, .s4};
65//! ----------------------------------------------------
66//! // Alternate floating point type:
67//! // Single bit:
68//! mma.sync.aligned.shape.row.col.s32.b1.b1.s32.bitOp.popc d, a, b, c;
69//! .bitOp = {.xor, .and};
70//! .shape = {.m8n8k128, .m16n8k128, .m16n8k256};
71
72#![allow(unused)]
73use crate::r#type::common::*;
74
75pub mod section_0 {
76    use crate::Spanned;
77    use crate::parser::Span;
78    use crate::r#type::common::*;
79
80    use serde::Serialize;
81
82    #[derive(Debug, Clone, PartialEq, Serialize)]
83    pub enum Alayout {
84        Row, // .row
85        Col, // .col
86    }
87
88    #[derive(Debug, Clone, PartialEq, Serialize)]
89    pub enum Blayout {
90        Row, // .row
91        Col, // .col
92    }
93
94    #[derive(Debug, Clone, PartialEq, Serialize)]
95    pub enum Dtype {
96        F16, // .f16
97        F32, // .f32
98    }
99
100    #[derive(Debug, Clone, PartialEq, Serialize)]
101    pub enum Ctype {
102        F16, // .f16
103        F32, // .f32
104    }
105
106    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
107    pub struct MmaSyncAlignedM8n8k4AlayoutBlayoutDtypeF16F16Ctype {
108        pub sync: (),          // .sync
109        pub aligned: (),       // .aligned
110        pub m8n8k4: (),        // .m8n8k4
111        pub alayout: Alayout,  // .alayout
112        pub blayout: Blayout,  // .blayout
113        pub dtype: Dtype,      // .dtype
114        pub f16: (),           // .f16
115        pub f162: (),          // .f16
116        pub ctype: Ctype,      // .ctype
117        pub d: GeneralOperand, // d
118        pub a: GeneralOperand, // a
119        pub b: GeneralOperand, // b
120        pub c: GeneralOperand, // c
121        pub span: Span,
122    }
123
124    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
125    pub struct MmaSyncAlignedM16n8k8RowColDtypeF16F16Ctype {
126        pub sync: (),          // .sync
127        pub aligned: (),       // .aligned
128        pub m16n8k8: (),       // .m16n8k8
129        pub row: (),           // .row
130        pub col: (),           // .col
131        pub dtype: Dtype,      // .dtype
132        pub f16: (),           // .f16
133        pub f162: (),          // .f16
134        pub ctype: Ctype,      // .ctype
135        pub d: GeneralOperand, // d
136        pub a: GeneralOperand, // a
137        pub b: GeneralOperand, // b
138        pub c: GeneralOperand, // c
139        pub span: Span,
140    }
141
142    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
143    pub struct MmaSyncAlignedM16n8k16RowColDtypeF16F16Ctype {
144        pub sync: (),          // .sync
145        pub aligned: (),       // .aligned
146        pub m16n8k16: (),      // .m16n8k16
147        pub row: (),           // .row
148        pub col: (),           // .col
149        pub dtype: Dtype,      // .dtype
150        pub f16: (),           // .f16
151        pub f162: (),          // .f16
152        pub ctype: Ctype,      // .ctype
153        pub d: GeneralOperand, // d
154        pub a: GeneralOperand, // a
155        pub b: GeneralOperand, // b
156        pub c: GeneralOperand, // c
157        pub span: Span,
158    }
159}
160
161pub mod section_1 {
162    use crate::Spanned;
163    use crate::parser::Span;
164    use crate::r#type::common::*;
165
166    use serde::Serialize;
167
168    #[derive(Debug, Clone, PartialEq, Serialize)]
169    pub enum Atype {
170        Bf16, // .bf16
171        Tf32, // .tf32
172    }
173
174    #[derive(Debug, Clone, PartialEq, Serialize)]
175    pub enum Btype {
176        Bf16, // .bf16
177        Tf32, // .tf32
178    }
179
180    #[derive(Debug, Clone, PartialEq, Serialize)]
181    pub enum Shape {
182        M16n8k16, // .m16n8k16
183        M16n8k32, // .m16n8k32
184    }
185
186    #[derive(Debug, Clone, PartialEq, Serialize)]
187    pub enum Dtype {
188        F16, // .f16
189        F32, // .f32
190    }
191
192    #[derive(Debug, Clone, PartialEq, Serialize)]
193    pub enum F8type {
194        E4m3, // .e4m3
195        E5m2, // .e5m2
196    }
197
198    #[derive(Debug, Clone, PartialEq, Serialize)]
199    pub enum Ctype {
200        F16, // .f16
201        F32, // .f32
202    }
203
204    #[derive(Debug, Clone, PartialEq, Serialize)]
205    pub enum Kind {
206        KindF8f6f4, // .kind::f8f6f4
207    }
208
209    #[derive(Debug, Clone, PartialEq, Serialize)]
210    pub enum F8f6f4type {
211        E4m3, // .e4m3
212        E5m2, // .e5m2
213        E3m2, // .e3m2
214        E2m3, // .e2m3
215        E2m1, // .e2m1
216    }
217
218    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
219    pub struct MmaSyncAlignedM16n8k4RowColF32Tf32Tf32F32 {
220        pub sync: (),          // .sync
221        pub aligned: (),       // .aligned
222        pub m16n8k4: (),       // .m16n8k4
223        pub row: (),           // .row
224        pub col: (),           // .col
225        pub f32: (),           // .f32
226        pub tf32: (),          // .tf32
227        pub tf322: (),         // .tf32
228        pub f322: (),          // .f32
229        pub d: GeneralOperand, // d
230        pub a: GeneralOperand, // a
231        pub b: GeneralOperand, // b
232        pub c: GeneralOperand, // c
233        pub span: Span,
234    }
235
236    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
237    pub struct MmaSyncAlignedM16n8k8RowColF32AtypeBtypeF32 {
238        pub sync: (),          // .sync
239        pub aligned: (),       // .aligned
240        pub m16n8k8: (),       // .m16n8k8
241        pub row: (),           // .row
242        pub col: (),           // .col
243        pub f32: (),           // .f32
244        pub atype: Atype,      // .atype
245        pub btype: Btype,      // .btype
246        pub f322: (),          // .f32
247        pub d: GeneralOperand, // d
248        pub a: GeneralOperand, // a
249        pub b: GeneralOperand, // b
250        pub c: GeneralOperand, // c
251        pub span: Span,
252    }
253
254    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
255    pub struct MmaSyncAlignedM16n8k16RowColF32Bf16Bf16F32 {
256        pub sync: (),          // .sync
257        pub aligned: (),       // .aligned
258        pub m16n8k16: (),      // .m16n8k16
259        pub row: (),           // .row
260        pub col: (),           // .col
261        pub f32: (),           // .f32
262        pub bf16: (),          // .bf16
263        pub bf162: (),         // .bf16
264        pub f322: (),          // .f32
265        pub d: GeneralOperand, // d
266        pub a: GeneralOperand, // a
267        pub b: GeneralOperand, // b
268        pub c: GeneralOperand, // c
269        pub span: Span,
270    }
271
272    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
273    pub struct MmaSyncAlignedShapeRowColDtypeF8typeF8typeCtype {
274        pub sync: (),          // .sync
275        pub aligned: (),       // .aligned
276        pub shape: Shape,      // .shape
277        pub row: (),           // .row
278        pub col: (),           // .col
279        pub dtype: Dtype,      // .dtype
280        pub f8type: F8type,    // .f8type
281        pub f8type1: F8type,   // .f8type
282        pub ctype: Ctype,      // .ctype
283        pub d: GeneralOperand, // d
284        pub a: GeneralOperand, // a
285        pub b: GeneralOperand, // b
286        pub c: GeneralOperand, // c
287        pub span: Span,
288    }
289
290    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
291    pub struct MmaSyncAlignedM16n8k32RowColKindDtypeF8f6f4typeF8f6f4typeCtype {
292        pub sync: (),                // .sync
293        pub aligned: (),             // .aligned
294        pub m16n8k32: (),            // .m16n8k32
295        pub row: (),                 // .row
296        pub col: (),                 // .col
297        pub kind: Kind,              // .kind
298        pub dtype: Dtype,            // .dtype
299        pub f8f6f4type: F8f6f4type,  // .f8f6f4type
300        pub f8f6f4type1: F8f6f4type, // .f8f6f4type
301        pub ctype: Ctype,            // .ctype
302        pub d: GeneralOperand,       // d
303        pub a: GeneralOperand,       // a
304        pub b: GeneralOperand,       // b
305        pub c: GeneralOperand,       // c
306        pub span: Span,
307    }
308}
309
310pub mod section_2 {
311    use crate::Spanned;
312    use crate::parser::Span;
313    use crate::r#type::common::*;
314
315    use serde::Serialize;
316
317    #[derive(Debug, Clone, PartialEq, Serialize)]
318    pub enum Kind {
319        KindMxf4, // .kind::mxf4
320    }
321
322    #[derive(Debug, Clone, PartialEq, Serialize)]
323    pub enum ScaleVecSize {
324        ScaleVec2x, // .scale_vec::2X
325    }
326
327    #[derive(Debug, Clone, PartialEq, Serialize)]
328    pub enum Stype {
329        Ue8m0, // .ue8m0
330    }
331
332    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
333    pub struct MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype {
334        pub sync: (),                             // .sync
335        pub aligned: (),                          // .aligned
336        pub m16n8k64: (),                         // .m16n8k64
337        pub row: (),                              // .row
338        pub col: (),                              // .col
339        pub kind: Kind,                           // .kind
340        pub block_scale: (),                      // .block_scale
341        pub scale_vec_size: Option<ScaleVecSize>, // {.scale_vec_size}
342        pub f32: (),                              // .f32
343        pub e2m1: (),                             // .e2m1
344        pub e2m12: (),                            // .e2m1
345        pub f322: (),                             // .f32
346        pub stype: Stype,                         // .stype
347        pub d: GeneralOperand,                    // d
348        pub a: GeneralOperand,                    // a
349        pub b: GeneralOperand,                    // b
350        pub c: GeneralOperand,                    // c
351        pub scale_a_data: GeneralOperand,         // scale-a-data
352        pub byte_id_a: VectorOperand,             // {byte-id-a, thread-id-a}
353        pub scale_b_data: GeneralOperand,         // scale-b-data
354        pub byte_id_b: VectorOperand,             // {byte-id-b, thread-id-b}
355        pub span: Span,
356    }
357}
358
359pub mod section_3 {
360    use crate::Spanned;
361    use crate::parser::Span;
362    use crate::r#type::common::*;
363
364    use serde::Serialize;
365
366    #[derive(Debug, Clone, PartialEq, Serialize)]
367    pub enum Kind {
368        KindMxf4nvf4, // .kind::mxf4nvf4
369    }
370
371    #[derive(Debug, Clone, PartialEq, Serialize)]
372    pub enum ScaleVecSize {
373        ScaleVec2x, // .scale_vec::2X
374        ScaleVec4x, // .scale_vec::4X
375    }
376
377    #[derive(Debug, Clone, PartialEq, Serialize)]
378    pub enum Stype {
379        Ue8m0, // .ue8m0
380        Ue4m3, // .ue4m3
381    }
382
383    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
384    pub struct MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype1 {
385        pub sync: (),                     // .sync
386        pub aligned: (),                  // .aligned
387        pub m16n8k64: (),                 // .m16n8k64
388        pub row: (),                      // .row
389        pub col: (),                      // .col
390        pub kind: Kind,                   // .kind
391        pub block_scale: (),              // .block_scale
392        pub scale_vec_size: ScaleVecSize, // .scale_vec_size
393        pub f32: (),                      // .f32
394        pub e2m1: (),                     // .e2m1
395        pub e2m12: (),                    // .e2m1
396        pub f322: (),                     // .f32
397        pub stype: Stype,                 // .stype
398        pub d: GeneralOperand,            // d
399        pub a: GeneralOperand,            // a
400        pub b: GeneralOperand,            // b
401        pub c: GeneralOperand,            // c
402        pub scale_a_data: GeneralOperand, // scale-a-data
403        pub byte_id_a: VectorOperand,     // {byte-id-a, thread-id-a}
404        pub scale_b_data: GeneralOperand, // scale-b-data
405        pub byte_id_b: VectorOperand,     // {byte-id-b, thread-id-b}
406        pub span: Span,
407    }
408}
409
410pub mod section_4 {
411    use crate::Spanned;
412    use crate::parser::Span;
413    use crate::r#type::common::*;
414
415    use serde::Serialize;
416
417    #[derive(Debug, Clone, PartialEq, Serialize)]
418    pub enum Kind {
419        KindMxf8f6f4, // .kind::mxf8f6f4
420    }
421
422    #[derive(Debug, Clone, PartialEq, Serialize)]
423    pub enum ScaleVecSize {
424        ScaleVec1x, // .scale_vec::1X
425    }
426
427    #[derive(Debug, Clone, PartialEq, Serialize)]
428    pub enum F8f6f4type {
429        E4m3, // .e4m3
430        E5m2, // .e5m2
431        E3m2, // .e3m2
432        E2m3, // .e2m3
433        E2m1, // .e2m1
434    }
435
436    #[derive(Debug, Clone, PartialEq, Serialize)]
437    pub enum Stype {
438        Ue8m0, // .ue8m0
439    }
440
441    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
442    pub struct MmaSyncAlignedM16n8k32RowColKindBlockScaleScaleVecSizeF32F8f6f4typeF8f6f4typeF32Stype {
443        pub sync: (),                             // .sync
444        pub aligned: (),                          // .aligned
445        pub m16n8k32: (),                         // .m16n8k32
446        pub row: (),                              // .row
447        pub col: (),                              // .col
448        pub kind: Kind,                           // .kind
449        pub block_scale: (),                      // .block_scale
450        pub scale_vec_size: Option<ScaleVecSize>, // {.scale_vec_size}
451        pub f32: (),                              // .f32
452        pub f8f6f4type: F8f6f4type,               // .f8f6f4type
453        pub f8f6f4type1: F8f6f4type,              // .f8f6f4type
454        pub f322: (),                             // .f32
455        pub stype: Stype,                         // .stype
456        pub d: GeneralOperand,                    // d
457        pub a: GeneralOperand,                    // a
458        pub b: GeneralOperand,                    // b
459        pub c: GeneralOperand,                    // c
460        pub scale_a_data: GeneralOperand,         // scale-a-data
461        pub byte_id_a: VectorOperand,             // {byte-id-a, thread-id-a}
462        pub scale_b_data: GeneralOperand,         // scale-b-data
463        pub byte_id_b: VectorOperand,             // {byte-id-b, thread-id-b}
464        pub span: Span,
465    }
466}
467
468pub mod section_5 {
469    use crate::Spanned;
470    use crate::parser::Span;
471    use crate::r#type::common::*;
472
473    use serde::Serialize;
474
475    #[derive(Debug, Clone, PartialEq, Serialize)]
476    pub enum Shape {
477        M16n8k16, // .m16n8k16
478        M16n8k4,  // .m16n8k4
479        M16n8k8,  // .m16n8k8
480        M8n84,    // .m8n84
481    }
482
483    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
484    pub struct MmaSyncAlignedShapeRowColF64F64F64F64 {
485        pub sync: (),          // .sync
486        pub aligned: (),       // .aligned
487        pub shape: Shape,      // .shape
488        pub row: (),           // .row
489        pub col: (),           // .col
490        pub f64: (),           // .f64
491        pub f642: (),          // .f64
492        pub f644: (),          // .f64
493        pub f646: (),          // .f64
494        pub d: GeneralOperand, // d
495        pub a: GeneralOperand, // a
496        pub b: GeneralOperand, // b
497        pub c: GeneralOperand, // c
498        pub span: Span,
499    }
500}
501
502pub mod section_6 {
503    use crate::Spanned;
504    use crate::parser::Span;
505    use crate::r#type::common::*;
506
507    use serde::Serialize;
508
509    #[derive(Debug, Clone, PartialEq, Serialize)]
510    pub enum Shape {
511        M16n8k16, // .m16n8k16
512        M16n8k32, // .m16n8k32
513        M8n8k16,  // .m8n8k16
514    }
515
516    #[derive(Debug, Clone, PartialEq, Serialize)]
517    pub enum Atype {
518        U8, // .u8
519        S8, // .s8
520    }
521
522    #[derive(Debug, Clone, PartialEq, Serialize)]
523    pub enum Btype {
524        U8, // .u8
525        S8, // .s8
526    }
527
528    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
529    pub struct MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS32 {
530        pub sync: (),          // .sync
531        pub aligned: (),       // .aligned
532        pub shape: Shape,      // .shape
533        pub row: (),           // .row
534        pub col: (),           // .col
535        pub satfinite: bool,   // {.satfinite}
536        pub s32: (),           // .s32
537        pub atype: Atype,      // .atype
538        pub btype: Btype,      // .btype
539        pub s322: (),          // .s32
540        pub d: GeneralOperand, // d
541        pub a: GeneralOperand, // a
542        pub b: GeneralOperand, // b
543        pub c: GeneralOperand, // c
544        pub span: Span,
545    }
546}
547
548pub mod section_7 {
549    use crate::Spanned;
550    use crate::parser::Span;
551    use crate::r#type::common::*;
552
553    use serde::Serialize;
554
555    #[derive(Debug, Clone, PartialEq, Serialize)]
556    pub enum Shape {
557        M16n8k32, // .m16n8k32
558        M16n8k64, // .m16n8k64
559        M8n8k32,  // .m8n8k32
560    }
561
562    #[derive(Debug, Clone, PartialEq, Serialize)]
563    pub enum Atype {
564        U4, // .u4
565        S4, // .s4
566    }
567
568    #[derive(Debug, Clone, PartialEq, Serialize)]
569    pub enum Btype {
570        U4, // .u4
571        S4, // .s4
572    }
573
574    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
575    pub struct MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS321 {
576        pub sync: (),          // .sync
577        pub aligned: (),       // .aligned
578        pub shape: Shape,      // .shape
579        pub row: (),           // .row
580        pub col: (),           // .col
581        pub satfinite: bool,   // {.satfinite}
582        pub s32: (),           // .s32
583        pub atype: Atype,      // .atype
584        pub btype: Btype,      // .btype
585        pub s322: (),          // .s32
586        pub d: GeneralOperand, // d
587        pub a: GeneralOperand, // a
588        pub b: GeneralOperand, // b
589        pub c: GeneralOperand, // c
590        pub span: Span,
591    }
592}
593
594pub mod section_8 {
595    use crate::Spanned;
596    use crate::parser::Span;
597    use crate::r#type::common::*;
598
599    use serde::Serialize;
600
601    #[derive(Debug, Clone, PartialEq, Serialize)]
602    pub enum Shape {
603        M16n8k128, // .m16n8k128
604        M16n8k256, // .m16n8k256
605        M8n8k128,  // .m8n8k128
606    }
607
608    #[derive(Debug, Clone, PartialEq, Serialize)]
609    pub enum Bitop {
610        Xor, // .xor
611        And, // .and
612    }
613
614    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
615    pub struct MmaSyncAlignedShapeRowColS32B1B1S32BitopPopc {
616        pub sync: (),          // .sync
617        pub aligned: (),       // .aligned
618        pub shape: Shape,      // .shape
619        pub row: (),           // .row
620        pub col: (),           // .col
621        pub s32: (),           // .s32
622        pub b1: (),            // .b1
623        pub b12: (),           // .b1
624        pub s322: (),          // .s32
625        pub bitop: Bitop,      // .bitOp
626        pub popc: (),          // .popc
627        pub d: GeneralOperand, // d
628        pub a: GeneralOperand, // a
629        pub b: GeneralOperand, // b
630        pub c: GeneralOperand, // c
631        pub span: Span,
632    }
633}
634
635// Re-export types with section suffixes to avoid naming conflicts
636// e.g., Type0 for section_0::Type, Type1 for section_1::Type
637pub use section_0::Alayout as Alayout0;
638pub use section_0::Blayout as Blayout0;
639pub use section_0::Ctype as Ctype0;
640pub use section_0::Dtype as Dtype0;
641pub use section_0::MmaSyncAlignedM8n8k4AlayoutBlayoutDtypeF16F16Ctype;
642pub use section_0::MmaSyncAlignedM16n8k8RowColDtypeF16F16Ctype;
643pub use section_0::MmaSyncAlignedM16n8k16RowColDtypeF16F16Ctype;
644pub use section_1::Atype as Atype1;
645pub use section_1::Btype as Btype1;
646pub use section_1::Ctype as Ctype1;
647pub use section_1::Dtype as Dtype1;
648pub use section_1::F8f6f4type as F8f6f4type1;
649pub use section_1::F8type as F8type1;
650pub use section_1::Kind as Kind1;
651pub use section_1::MmaSyncAlignedM16n8k4RowColF32Tf32Tf32F32;
652pub use section_1::MmaSyncAlignedM16n8k8RowColF32AtypeBtypeF32;
653pub use section_1::MmaSyncAlignedM16n8k16RowColF32Bf16Bf16F32;
654pub use section_1::MmaSyncAlignedM16n8k32RowColKindDtypeF8f6f4typeF8f6f4typeCtype;
655pub use section_1::MmaSyncAlignedShapeRowColDtypeF8typeF8typeCtype;
656pub use section_1::Shape as Shape1;
657pub use section_2::Kind as Kind2;
658pub use section_2::MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype;
659pub use section_2::ScaleVecSize as ScaleVecSize2;
660pub use section_2::Stype as Stype2;
661pub use section_3::Kind as Kind3;
662pub use section_3::MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype1;
663pub use section_3::ScaleVecSize as ScaleVecSize3;
664pub use section_3::Stype as Stype3;
665pub use section_4::F8f6f4type as F8f6f4type4;
666pub use section_4::Kind as Kind4;
667pub use section_4::MmaSyncAlignedM16n8k32RowColKindBlockScaleScaleVecSizeF32F8f6f4typeF8f6f4typeF32Stype;
668pub use section_4::ScaleVecSize as ScaleVecSize4;
669pub use section_4::Stype as Stype4;
670pub use section_5::MmaSyncAlignedShapeRowColF64F64F64F64;
671pub use section_5::Shape as Shape5;
672pub use section_6::Atype as Atype6;
673pub use section_6::Btype as Btype6;
674pub use section_6::MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS32;
675pub use section_6::Shape as Shape6;
676pub use section_7::Atype as Atype7;
677pub use section_7::Btype as Btype7;
678pub use section_7::MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS321;
679pub use section_7::Shape as Shape7;
680pub use section_8::Bitop as Bitop8;
681pub use section_8::MmaSyncAlignedShapeRowColS32B1B1S32BitopPopc;
682pub use section_8::Shape as Shape8;
ptx_parser/type/instruction/mma.rs

ptx_parser/type/instruction/
mma.rs
