ptx_parser/unparser/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)]
73
74use crate::lexer::PtxToken;
75use crate::unparser::{PtxUnparser, common::*};
76
77pub mod section_0 {
78    use super::*;
79    use crate::r#type::instruction::mma::section_0::*;
80
81    impl PtxUnparser for MmaSyncAlignedM8n8k4AlayoutBlayoutDtypeF16F16Ctype {
82        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
83            self.unparse_tokens_mode(tokens, false);
84        }
85        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
86            push_opcode(tokens, "mma");
87            push_directive(tokens, "sync");
88            push_directive(tokens, "aligned");
89            push_directive(tokens, "m8n8k4");
90            match &self.alayout {
91                Alayout::Row => {
92                    push_directive(tokens, "row");
93                }
94                Alayout::Col => {
95                    push_directive(tokens, "col");
96                }
97            }
98            match &self.blayout {
99                Blayout::Row => {
100                    push_directive(tokens, "row");
101                }
102                Blayout::Col => {
103                    push_directive(tokens, "col");
104                }
105            }
106            match &self.dtype {
107                Dtype::F16 => {
108                    push_directive(tokens, "f16");
109                }
110                Dtype::F32 => {
111                    push_directive(tokens, "f32");
112                }
113            }
114            push_directive(tokens, "f16");
115            push_directive(tokens, "f16");
116            match &self.ctype {
117                Ctype::F16 => {
118                    push_directive(tokens, "f16");
119                }
120                Ctype::F32 => {
121                    push_directive(tokens, "f32");
122                }
123            }
124            if spaced {
125                tokens.push(PtxToken::Space);
126            }
127            self.d.unparse_tokens_mode(tokens, spaced);
128            tokens.push(PtxToken::Comma);
129            if spaced {
130                tokens.push(PtxToken::Space);
131            }
132            self.a.unparse_tokens_mode(tokens, spaced);
133            tokens.push(PtxToken::Comma);
134            if spaced {
135                tokens.push(PtxToken::Space);
136            }
137            self.b.unparse_tokens_mode(tokens, spaced);
138            tokens.push(PtxToken::Comma);
139            if spaced {
140                tokens.push(PtxToken::Space);
141            }
142            self.c.unparse_tokens_mode(tokens, spaced);
143            tokens.push(PtxToken::Semicolon);
144            if spaced {
145                tokens.push(PtxToken::Newline);
146            }
147        }
148    }
149
150    impl PtxUnparser for MmaSyncAlignedM16n8k8RowColDtypeF16F16Ctype {
151        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
152            self.unparse_tokens_mode(tokens, false);
153        }
154        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
155            push_opcode(tokens, "mma");
156            push_directive(tokens, "sync");
157            push_directive(tokens, "aligned");
158            push_directive(tokens, "m16n8k8");
159            push_directive(tokens, "row");
160            push_directive(tokens, "col");
161            match &self.dtype {
162                Dtype::F16 => {
163                    push_directive(tokens, "f16");
164                }
165                Dtype::F32 => {
166                    push_directive(tokens, "f32");
167                }
168            }
169            push_directive(tokens, "f16");
170            push_directive(tokens, "f16");
171            match &self.ctype {
172                Ctype::F16 => {
173                    push_directive(tokens, "f16");
174                }
175                Ctype::F32 => {
176                    push_directive(tokens, "f32");
177                }
178            }
179            if spaced {
180                tokens.push(PtxToken::Space);
181            }
182            self.d.unparse_tokens_mode(tokens, spaced);
183            tokens.push(PtxToken::Comma);
184            if spaced {
185                tokens.push(PtxToken::Space);
186            }
187            self.a.unparse_tokens_mode(tokens, spaced);
188            tokens.push(PtxToken::Comma);
189            if spaced {
190                tokens.push(PtxToken::Space);
191            }
192            self.b.unparse_tokens_mode(tokens, spaced);
193            tokens.push(PtxToken::Comma);
194            if spaced {
195                tokens.push(PtxToken::Space);
196            }
197            self.c.unparse_tokens_mode(tokens, spaced);
198            tokens.push(PtxToken::Semicolon);
199            if spaced {
200                tokens.push(PtxToken::Newline);
201            }
202        }
203    }
204
205    impl PtxUnparser for MmaSyncAlignedM16n8k16RowColDtypeF16F16Ctype {
206        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
207            self.unparse_tokens_mode(tokens, false);
208        }
209        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
210            push_opcode(tokens, "mma");
211            push_directive(tokens, "sync");
212            push_directive(tokens, "aligned");
213            push_directive(tokens, "m16n8k16");
214            push_directive(tokens, "row");
215            push_directive(tokens, "col");
216            match &self.dtype {
217                Dtype::F16 => {
218                    push_directive(tokens, "f16");
219                }
220                Dtype::F32 => {
221                    push_directive(tokens, "f32");
222                }
223            }
224            push_directive(tokens, "f16");
225            push_directive(tokens, "f16");
226            match &self.ctype {
227                Ctype::F16 => {
228                    push_directive(tokens, "f16");
229                }
230                Ctype::F32 => {
231                    push_directive(tokens, "f32");
232                }
233            }
234            if spaced {
235                tokens.push(PtxToken::Space);
236            }
237            self.d.unparse_tokens_mode(tokens, spaced);
238            tokens.push(PtxToken::Comma);
239            if spaced {
240                tokens.push(PtxToken::Space);
241            }
242            self.a.unparse_tokens_mode(tokens, spaced);
243            tokens.push(PtxToken::Comma);
244            if spaced {
245                tokens.push(PtxToken::Space);
246            }
247            self.b.unparse_tokens_mode(tokens, spaced);
248            tokens.push(PtxToken::Comma);
249            if spaced {
250                tokens.push(PtxToken::Space);
251            }
252            self.c.unparse_tokens_mode(tokens, spaced);
253            tokens.push(PtxToken::Semicolon);
254            if spaced {
255                tokens.push(PtxToken::Newline);
256            }
257        }
258    }
259}
260
261pub mod section_1 {
262    use super::*;
263    use crate::r#type::instruction::mma::section_1::*;
264
265    impl PtxUnparser for MmaSyncAlignedM16n8k4RowColF32Tf32Tf32F32 {
266        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
267            self.unparse_tokens_mode(tokens, false);
268        }
269        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
270            push_opcode(tokens, "mma");
271            push_directive(tokens, "sync");
272            push_directive(tokens, "aligned");
273            push_directive(tokens, "m16n8k4");
274            push_directive(tokens, "row");
275            push_directive(tokens, "col");
276            push_directive(tokens, "f32");
277            push_directive(tokens, "tf32");
278            push_directive(tokens, "tf32");
279            push_directive(tokens, "f32");
280            if spaced {
281                tokens.push(PtxToken::Space);
282            }
283            self.d.unparse_tokens_mode(tokens, spaced);
284            tokens.push(PtxToken::Comma);
285            if spaced {
286                tokens.push(PtxToken::Space);
287            }
288            self.a.unparse_tokens_mode(tokens, spaced);
289            tokens.push(PtxToken::Comma);
290            if spaced {
291                tokens.push(PtxToken::Space);
292            }
293            self.b.unparse_tokens_mode(tokens, spaced);
294            tokens.push(PtxToken::Comma);
295            if spaced {
296                tokens.push(PtxToken::Space);
297            }
298            self.c.unparse_tokens_mode(tokens, spaced);
299            tokens.push(PtxToken::Semicolon);
300            if spaced {
301                tokens.push(PtxToken::Newline);
302            }
303        }
304    }
305
306    impl PtxUnparser for MmaSyncAlignedM16n8k8RowColF32AtypeBtypeF32 {
307        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
308            self.unparse_tokens_mode(tokens, false);
309        }
310        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
311            push_opcode(tokens, "mma");
312            push_directive(tokens, "sync");
313            push_directive(tokens, "aligned");
314            push_directive(tokens, "m16n8k8");
315            push_directive(tokens, "row");
316            push_directive(tokens, "col");
317            push_directive(tokens, "f32");
318            match &self.atype {
319                Atype::Bf16 => {
320                    push_directive(tokens, "bf16");
321                }
322                Atype::Tf32 => {
323                    push_directive(tokens, "tf32");
324                }
325            }
326            match &self.btype {
327                Btype::Bf16 => {
328                    push_directive(tokens, "bf16");
329                }
330                Btype::Tf32 => {
331                    push_directive(tokens, "tf32");
332                }
333            }
334            push_directive(tokens, "f32");
335            if spaced {
336                tokens.push(PtxToken::Space);
337            }
338            self.d.unparse_tokens_mode(tokens, spaced);
339            tokens.push(PtxToken::Comma);
340            if spaced {
341                tokens.push(PtxToken::Space);
342            }
343            self.a.unparse_tokens_mode(tokens, spaced);
344            tokens.push(PtxToken::Comma);
345            if spaced {
346                tokens.push(PtxToken::Space);
347            }
348            self.b.unparse_tokens_mode(tokens, spaced);
349            tokens.push(PtxToken::Comma);
350            if spaced {
351                tokens.push(PtxToken::Space);
352            }
353            self.c.unparse_tokens_mode(tokens, spaced);
354            tokens.push(PtxToken::Semicolon);
355            if spaced {
356                tokens.push(PtxToken::Newline);
357            }
358        }
359    }
360
361    impl PtxUnparser for MmaSyncAlignedM16n8k16RowColF32Bf16Bf16F32 {
362        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
363            self.unparse_tokens_mode(tokens, false);
364        }
365        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
366            push_opcode(tokens, "mma");
367            push_directive(tokens, "sync");
368            push_directive(tokens, "aligned");
369            push_directive(tokens, "m16n8k16");
370            push_directive(tokens, "row");
371            push_directive(tokens, "col");
372            push_directive(tokens, "f32");
373            push_directive(tokens, "bf16");
374            push_directive(tokens, "bf16");
375            push_directive(tokens, "f32");
376            if spaced {
377                tokens.push(PtxToken::Space);
378            }
379            self.d.unparse_tokens_mode(tokens, spaced);
380            tokens.push(PtxToken::Comma);
381            if spaced {
382                tokens.push(PtxToken::Space);
383            }
384            self.a.unparse_tokens_mode(tokens, spaced);
385            tokens.push(PtxToken::Comma);
386            if spaced {
387                tokens.push(PtxToken::Space);
388            }
389            self.b.unparse_tokens_mode(tokens, spaced);
390            tokens.push(PtxToken::Comma);
391            if spaced {
392                tokens.push(PtxToken::Space);
393            }
394            self.c.unparse_tokens_mode(tokens, spaced);
395            tokens.push(PtxToken::Semicolon);
396            if spaced {
397                tokens.push(PtxToken::Newline);
398            }
399        }
400    }
401
402    impl PtxUnparser for MmaSyncAlignedShapeRowColDtypeF8typeF8typeCtype {
403        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
404            self.unparse_tokens_mode(tokens, false);
405        }
406        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
407            push_opcode(tokens, "mma");
408            push_directive(tokens, "sync");
409            push_directive(tokens, "aligned");
410            match &self.shape {
411                Shape::M16n8k16 => {
412                    push_directive(tokens, "m16n8k16");
413                }
414                Shape::M16n8k32 => {
415                    push_directive(tokens, "m16n8k32");
416                }
417            }
418            push_directive(tokens, "row");
419            push_directive(tokens, "col");
420            match &self.dtype {
421                Dtype::F16 => {
422                    push_directive(tokens, "f16");
423                }
424                Dtype::F32 => {
425                    push_directive(tokens, "f32");
426                }
427            }
428            match &self.f8type {
429                F8type::E4m3 => {
430                    push_directive(tokens, "e4m3");
431                }
432                F8type::E5m2 => {
433                    push_directive(tokens, "e5m2");
434                }
435            }
436            match &self.f8type1 {
437                F8type::E4m3 => {
438                    push_directive(tokens, "e4m3");
439                }
440                F8type::E5m2 => {
441                    push_directive(tokens, "e5m2");
442                }
443            }
444            match &self.ctype {
445                Ctype::F16 => {
446                    push_directive(tokens, "f16");
447                }
448                Ctype::F32 => {
449                    push_directive(tokens, "f32");
450                }
451            }
452            if spaced {
453                tokens.push(PtxToken::Space);
454            }
455            self.d.unparse_tokens_mode(tokens, spaced);
456            tokens.push(PtxToken::Comma);
457            if spaced {
458                tokens.push(PtxToken::Space);
459            }
460            self.a.unparse_tokens_mode(tokens, spaced);
461            tokens.push(PtxToken::Comma);
462            if spaced {
463                tokens.push(PtxToken::Space);
464            }
465            self.b.unparse_tokens_mode(tokens, spaced);
466            tokens.push(PtxToken::Comma);
467            if spaced {
468                tokens.push(PtxToken::Space);
469            }
470            self.c.unparse_tokens_mode(tokens, spaced);
471            tokens.push(PtxToken::Semicolon);
472            if spaced {
473                tokens.push(PtxToken::Newline);
474            }
475        }
476    }
477
478    impl PtxUnparser for MmaSyncAlignedM16n8k32RowColKindDtypeF8f6f4typeF8f6f4typeCtype {
479        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
480            self.unparse_tokens_mode(tokens, false);
481        }
482        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
483            push_opcode(tokens, "mma");
484            push_directive(tokens, "sync");
485            push_directive(tokens, "aligned");
486            push_directive(tokens, "m16n8k32");
487            push_directive(tokens, "row");
488            push_directive(tokens, "col");
489            match &self.kind {
490                Kind::KindF8f6f4 => {
491                    push_directive(tokens, "kind::f8f6f4");
492                }
493            }
494            match &self.dtype {
495                Dtype::F16 => {
496                    push_directive(tokens, "f16");
497                }
498                Dtype::F32 => {
499                    push_directive(tokens, "f32");
500                }
501            }
502            match &self.f8f6f4type {
503                F8f6f4type::E4m3 => {
504                    push_directive(tokens, "e4m3");
505                }
506                F8f6f4type::E5m2 => {
507                    push_directive(tokens, "e5m2");
508                }
509                F8f6f4type::E3m2 => {
510                    push_directive(tokens, "e3m2");
511                }
512                F8f6f4type::E2m3 => {
513                    push_directive(tokens, "e2m3");
514                }
515                F8f6f4type::E2m1 => {
516                    push_directive(tokens, "e2m1");
517                }
518            }
519            match &self.f8f6f4type1 {
520                F8f6f4type::E4m3 => {
521                    push_directive(tokens, "e4m3");
522                }
523                F8f6f4type::E5m2 => {
524                    push_directive(tokens, "e5m2");
525                }
526                F8f6f4type::E3m2 => {
527                    push_directive(tokens, "e3m2");
528                }
529                F8f6f4type::E2m3 => {
530                    push_directive(tokens, "e2m3");
531                }
532                F8f6f4type::E2m1 => {
533                    push_directive(tokens, "e2m1");
534                }
535            }
536            match &self.ctype {
537                Ctype::F16 => {
538                    push_directive(tokens, "f16");
539                }
540                Ctype::F32 => {
541                    push_directive(tokens, "f32");
542                }
543            }
544            if spaced {
545                tokens.push(PtxToken::Space);
546            }
547            self.d.unparse_tokens_mode(tokens, spaced);
548            tokens.push(PtxToken::Comma);
549            if spaced {
550                tokens.push(PtxToken::Space);
551            }
552            self.a.unparse_tokens_mode(tokens, spaced);
553            tokens.push(PtxToken::Comma);
554            if spaced {
555                tokens.push(PtxToken::Space);
556            }
557            self.b.unparse_tokens_mode(tokens, spaced);
558            tokens.push(PtxToken::Comma);
559            if spaced {
560                tokens.push(PtxToken::Space);
561            }
562            self.c.unparse_tokens_mode(tokens, spaced);
563            tokens.push(PtxToken::Semicolon);
564            if spaced {
565                tokens.push(PtxToken::Newline);
566            }
567        }
568    }
569}
570
571pub mod section_2 {
572    use super::*;
573    use crate::r#type::instruction::mma::section_2::*;
574
575    impl PtxUnparser for MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype {
576        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
577            self.unparse_tokens_mode(tokens, false);
578        }
579        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
580            push_opcode(tokens, "mma");
581            push_directive(tokens, "sync");
582            push_directive(tokens, "aligned");
583            push_directive(tokens, "m16n8k64");
584            push_directive(tokens, "row");
585            push_directive(tokens, "col");
586            match &self.kind {
587                Kind::KindMxf4 => {
588                    push_directive(tokens, "kind::mxf4");
589                }
590            }
591            push_directive(tokens, "block_scale");
592            if let Some(scale_vec_size_0) = self.scale_vec_size.as_ref() {
593                match scale_vec_size_0 {
594                    ScaleVecSize::ScaleVec2x => {
595                        push_directive(tokens, "scale_vec::2X");
596                    }
597                }
598            }
599            push_directive(tokens, "f32");
600            push_directive(tokens, "e2m1");
601            push_directive(tokens, "e2m1");
602            push_directive(tokens, "f32");
603            match &self.stype {
604                Stype::Ue8m0 => {
605                    push_directive(tokens, "ue8m0");
606                }
607            }
608            if spaced {
609                tokens.push(PtxToken::Space);
610            }
611            self.d.unparse_tokens_mode(tokens, spaced);
612            tokens.push(PtxToken::Comma);
613            if spaced {
614                tokens.push(PtxToken::Space);
615            }
616            self.a.unparse_tokens_mode(tokens, spaced);
617            tokens.push(PtxToken::Comma);
618            if spaced {
619                tokens.push(PtxToken::Space);
620            }
621            self.b.unparse_tokens_mode(tokens, spaced);
622            tokens.push(PtxToken::Comma);
623            if spaced {
624                tokens.push(PtxToken::Space);
625            }
626            self.c.unparse_tokens_mode(tokens, spaced);
627            tokens.push(PtxToken::Comma);
628            if spaced {
629                tokens.push(PtxToken::Space);
630            }
631            self.scale_a_data.unparse_tokens_mode(tokens, spaced);
632            tokens.push(PtxToken::Comma);
633            if spaced {
634                tokens.push(PtxToken::Space);
635            }
636            self.byte_id_a.unparse_tokens_mode(tokens, spaced);
637            tokens.push(PtxToken::Comma);
638            if spaced {
639                tokens.push(PtxToken::Space);
640            }
641            self.scale_b_data.unparse_tokens_mode(tokens, spaced);
642            tokens.push(PtxToken::Comma);
643            if spaced {
644                tokens.push(PtxToken::Space);
645            }
646            self.byte_id_b.unparse_tokens_mode(tokens, spaced);
647            tokens.push(PtxToken::Semicolon);
648            if spaced {
649                tokens.push(PtxToken::Newline);
650            }
651        }
652    }
653}
654
655pub mod section_3 {
656    use super::*;
657    use crate::r#type::instruction::mma::section_3::*;
658
659    impl PtxUnparser for MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype1 {
660        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
661            self.unparse_tokens_mode(tokens, false);
662        }
663        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
664            push_opcode(tokens, "mma");
665            push_directive(tokens, "sync");
666            push_directive(tokens, "aligned");
667            push_directive(tokens, "m16n8k64");
668            push_directive(tokens, "row");
669            push_directive(tokens, "col");
670            match &self.kind {
671                Kind::KindMxf4nvf4 => {
672                    push_directive(tokens, "kind::mxf4nvf4");
673                }
674            }
675            push_directive(tokens, "block_scale");
676            match &self.scale_vec_size {
677                ScaleVecSize::ScaleVec2x => {
678                    push_directive(tokens, "scale_vec::2X");
679                }
680                ScaleVecSize::ScaleVec4x => {
681                    push_directive(tokens, "scale_vec::4X");
682                }
683            }
684            push_directive(tokens, "f32");
685            push_directive(tokens, "e2m1");
686            push_directive(tokens, "e2m1");
687            push_directive(tokens, "f32");
688            match &self.stype {
689                Stype::Ue8m0 => {
690                    push_directive(tokens, "ue8m0");
691                }
692                Stype::Ue4m3 => {
693                    push_directive(tokens, "ue4m3");
694                }
695            }
696            if spaced {
697                tokens.push(PtxToken::Space);
698            }
699            self.d.unparse_tokens_mode(tokens, spaced);
700            tokens.push(PtxToken::Comma);
701            if spaced {
702                tokens.push(PtxToken::Space);
703            }
704            self.a.unparse_tokens_mode(tokens, spaced);
705            tokens.push(PtxToken::Comma);
706            if spaced {
707                tokens.push(PtxToken::Space);
708            }
709            self.b.unparse_tokens_mode(tokens, spaced);
710            tokens.push(PtxToken::Comma);
711            if spaced {
712                tokens.push(PtxToken::Space);
713            }
714            self.c.unparse_tokens_mode(tokens, spaced);
715            tokens.push(PtxToken::Comma);
716            if spaced {
717                tokens.push(PtxToken::Space);
718            }
719            self.scale_a_data.unparse_tokens_mode(tokens, spaced);
720            tokens.push(PtxToken::Comma);
721            if spaced {
722                tokens.push(PtxToken::Space);
723            }
724            self.byte_id_a.unparse_tokens_mode(tokens, spaced);
725            tokens.push(PtxToken::Comma);
726            if spaced {
727                tokens.push(PtxToken::Space);
728            }
729            self.scale_b_data.unparse_tokens_mode(tokens, spaced);
730            tokens.push(PtxToken::Comma);
731            if spaced {
732                tokens.push(PtxToken::Space);
733            }
734            self.byte_id_b.unparse_tokens_mode(tokens, spaced);
735            tokens.push(PtxToken::Semicolon);
736            if spaced {
737                tokens.push(PtxToken::Newline);
738            }
739        }
740    }
741}
742
743pub mod section_4 {
744    use super::*;
745    use crate::r#type::instruction::mma::section_4::*;
746
747    impl PtxUnparser
748        for MmaSyncAlignedM16n8k32RowColKindBlockScaleScaleVecSizeF32F8f6f4typeF8f6f4typeF32Stype
749    {
750        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
751            self.unparse_tokens_mode(tokens, false);
752        }
753        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
754            push_opcode(tokens, "mma");
755            push_directive(tokens, "sync");
756            push_directive(tokens, "aligned");
757            push_directive(tokens, "m16n8k32");
758            push_directive(tokens, "row");
759            push_directive(tokens, "col");
760            match &self.kind {
761                Kind::KindMxf8f6f4 => {
762                    push_directive(tokens, "kind::mxf8f6f4");
763                }
764            }
765            push_directive(tokens, "block_scale");
766            if let Some(scale_vec_size_1) = self.scale_vec_size.as_ref() {
767                match scale_vec_size_1 {
768                    ScaleVecSize::ScaleVec1x => {
769                        push_directive(tokens, "scale_vec::1X");
770                    }
771                }
772            }
773            push_directive(tokens, "f32");
774            match &self.f8f6f4type {
775                F8f6f4type::E4m3 => {
776                    push_directive(tokens, "e4m3");
777                }
778                F8f6f4type::E5m2 => {
779                    push_directive(tokens, "e5m2");
780                }
781                F8f6f4type::E3m2 => {
782                    push_directive(tokens, "e3m2");
783                }
784                F8f6f4type::E2m3 => {
785                    push_directive(tokens, "e2m3");
786                }
787                F8f6f4type::E2m1 => {
788                    push_directive(tokens, "e2m1");
789                }
790            }
791            match &self.f8f6f4type1 {
792                F8f6f4type::E4m3 => {
793                    push_directive(tokens, "e4m3");
794                }
795                F8f6f4type::E5m2 => {
796                    push_directive(tokens, "e5m2");
797                }
798                F8f6f4type::E3m2 => {
799                    push_directive(tokens, "e3m2");
800                }
801                F8f6f4type::E2m3 => {
802                    push_directive(tokens, "e2m3");
803                }
804                F8f6f4type::E2m1 => {
805                    push_directive(tokens, "e2m1");
806                }
807            }
808            push_directive(tokens, "f32");
809            match &self.stype {
810                Stype::Ue8m0 => {
811                    push_directive(tokens, "ue8m0");
812                }
813            }
814            if spaced {
815                tokens.push(PtxToken::Space);
816            }
817            self.d.unparse_tokens_mode(tokens, spaced);
818            tokens.push(PtxToken::Comma);
819            if spaced {
820                tokens.push(PtxToken::Space);
821            }
822            self.a.unparse_tokens_mode(tokens, spaced);
823            tokens.push(PtxToken::Comma);
824            if spaced {
825                tokens.push(PtxToken::Space);
826            }
827            self.b.unparse_tokens_mode(tokens, spaced);
828            tokens.push(PtxToken::Comma);
829            if spaced {
830                tokens.push(PtxToken::Space);
831            }
832            self.c.unparse_tokens_mode(tokens, spaced);
833            tokens.push(PtxToken::Comma);
834            if spaced {
835                tokens.push(PtxToken::Space);
836            }
837            self.scale_a_data.unparse_tokens_mode(tokens, spaced);
838            tokens.push(PtxToken::Comma);
839            if spaced {
840                tokens.push(PtxToken::Space);
841            }
842            self.byte_id_a.unparse_tokens_mode(tokens, spaced);
843            tokens.push(PtxToken::Comma);
844            if spaced {
845                tokens.push(PtxToken::Space);
846            }
847            self.scale_b_data.unparse_tokens_mode(tokens, spaced);
848            tokens.push(PtxToken::Comma);
849            if spaced {
850                tokens.push(PtxToken::Space);
851            }
852            self.byte_id_b.unparse_tokens_mode(tokens, spaced);
853            tokens.push(PtxToken::Semicolon);
854            if spaced {
855                tokens.push(PtxToken::Newline);
856            }
857        }
858    }
859}
860
861pub mod section_5 {
862    use super::*;
863    use crate::r#type::instruction::mma::section_5::*;
864
865    impl PtxUnparser for MmaSyncAlignedShapeRowColF64F64F64F64 {
866        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
867            self.unparse_tokens_mode(tokens, false);
868        }
869        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
870            push_opcode(tokens, "mma");
871            push_directive(tokens, "sync");
872            push_directive(tokens, "aligned");
873            match &self.shape {
874                Shape::M16n8k16 => {
875                    push_directive(tokens, "m16n8k16");
876                }
877                Shape::M16n8k4 => {
878                    push_directive(tokens, "m16n8k4");
879                }
880                Shape::M16n8k8 => {
881                    push_directive(tokens, "m16n8k8");
882                }
883                Shape::M8n84 => {
884                    push_directive(tokens, "m8n84");
885                }
886            }
887            push_directive(tokens, "row");
888            push_directive(tokens, "col");
889            push_directive(tokens, "f64");
890            push_directive(tokens, "f64");
891            push_directive(tokens, "f64");
892            push_directive(tokens, "f64");
893            if spaced {
894                tokens.push(PtxToken::Space);
895            }
896            self.d.unparse_tokens_mode(tokens, spaced);
897            tokens.push(PtxToken::Comma);
898            if spaced {
899                tokens.push(PtxToken::Space);
900            }
901            self.a.unparse_tokens_mode(tokens, spaced);
902            tokens.push(PtxToken::Comma);
903            if spaced {
904                tokens.push(PtxToken::Space);
905            }
906            self.b.unparse_tokens_mode(tokens, spaced);
907            tokens.push(PtxToken::Comma);
908            if spaced {
909                tokens.push(PtxToken::Space);
910            }
911            self.c.unparse_tokens_mode(tokens, spaced);
912            tokens.push(PtxToken::Semicolon);
913            if spaced {
914                tokens.push(PtxToken::Newline);
915            }
916        }
917    }
918}
919
920pub mod section_6 {
921    use super::*;
922    use crate::r#type::instruction::mma::section_6::*;
923
924    impl PtxUnparser for MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS32 {
925        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
926            self.unparse_tokens_mode(tokens, false);
927        }
928        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
929            push_opcode(tokens, "mma");
930            push_directive(tokens, "sync");
931            push_directive(tokens, "aligned");
932            match &self.shape {
933                Shape::M16n8k16 => {
934                    push_directive(tokens, "m16n8k16");
935                }
936                Shape::M16n8k32 => {
937                    push_directive(tokens, "m16n8k32");
938                }
939                Shape::M8n8k16 => {
940                    push_directive(tokens, "m8n8k16");
941                }
942            }
943            push_directive(tokens, "row");
944            push_directive(tokens, "col");
945            if self.satfinite {
946                push_directive(tokens, "satfinite");
947            }
948            push_directive(tokens, "s32");
949            match &self.atype {
950                Atype::U8 => {
951                    push_directive(tokens, "u8");
952                }
953                Atype::S8 => {
954                    push_directive(tokens, "s8");
955                }
956            }
957            match &self.btype {
958                Btype::U8 => {
959                    push_directive(tokens, "u8");
960                }
961                Btype::S8 => {
962                    push_directive(tokens, "s8");
963                }
964            }
965            push_directive(tokens, "s32");
966            if spaced {
967                tokens.push(PtxToken::Space);
968            }
969            self.d.unparse_tokens_mode(tokens, spaced);
970            tokens.push(PtxToken::Comma);
971            if spaced {
972                tokens.push(PtxToken::Space);
973            }
974            self.a.unparse_tokens_mode(tokens, spaced);
975            tokens.push(PtxToken::Comma);
976            if spaced {
977                tokens.push(PtxToken::Space);
978            }
979            self.b.unparse_tokens_mode(tokens, spaced);
980            tokens.push(PtxToken::Comma);
981            if spaced {
982                tokens.push(PtxToken::Space);
983            }
984            self.c.unparse_tokens_mode(tokens, spaced);
985            tokens.push(PtxToken::Semicolon);
986            if spaced {
987                tokens.push(PtxToken::Newline);
988            }
989        }
990    }
991}
992
993pub mod section_7 {
994    use super::*;
995    use crate::r#type::instruction::mma::section_7::*;
996
997    impl PtxUnparser for MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS321 {
998        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
999            self.unparse_tokens_mode(tokens, false);
1000        }
1001        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
1002            push_opcode(tokens, "mma");
1003            push_directive(tokens, "sync");
1004            push_directive(tokens, "aligned");
1005            match &self.shape {
1006                Shape::M16n8k32 => {
1007                    push_directive(tokens, "m16n8k32");
1008                }
1009                Shape::M16n8k64 => {
1010                    push_directive(tokens, "m16n8k64");
1011                }
1012                Shape::M8n8k32 => {
1013                    push_directive(tokens, "m8n8k32");
1014                }
1015            }
1016            push_directive(tokens, "row");
1017            push_directive(tokens, "col");
1018            if self.satfinite {
1019                push_directive(tokens, "satfinite");
1020            }
1021            push_directive(tokens, "s32");
1022            match &self.atype {
1023                Atype::U4 => {
1024                    push_directive(tokens, "u4");
1025                }
1026                Atype::S4 => {
1027                    push_directive(tokens, "s4");
1028                }
1029            }
1030            match &self.btype {
1031                Btype::U4 => {
1032                    push_directive(tokens, "u4");
1033                }
1034                Btype::S4 => {
1035                    push_directive(tokens, "s4");
1036                }
1037            }
1038            push_directive(tokens, "s32");
1039            if spaced {
1040                tokens.push(PtxToken::Space);
1041            }
1042            self.d.unparse_tokens_mode(tokens, spaced);
1043            tokens.push(PtxToken::Comma);
1044            if spaced {
1045                tokens.push(PtxToken::Space);
1046            }
1047            self.a.unparse_tokens_mode(tokens, spaced);
1048            tokens.push(PtxToken::Comma);
1049            if spaced {
1050                tokens.push(PtxToken::Space);
1051            }
1052            self.b.unparse_tokens_mode(tokens, spaced);
1053            tokens.push(PtxToken::Comma);
1054            if spaced {
1055                tokens.push(PtxToken::Space);
1056            }
1057            self.c.unparse_tokens_mode(tokens, spaced);
1058            tokens.push(PtxToken::Semicolon);
1059            if spaced {
1060                tokens.push(PtxToken::Newline);
1061            }
1062        }
1063    }
1064}
1065
1066pub mod section_8 {
1067    use super::*;
1068    use crate::r#type::instruction::mma::section_8::*;
1069
1070    impl PtxUnparser for MmaSyncAlignedShapeRowColS32B1B1S32BitopPopc {
1071        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
1072            self.unparse_tokens_mode(tokens, false);
1073        }
1074        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
1075            push_opcode(tokens, "mma");
1076            push_directive(tokens, "sync");
1077            push_directive(tokens, "aligned");
1078            match &self.shape {
1079                Shape::M16n8k128 => {
1080                    push_directive(tokens, "m16n8k128");
1081                }
1082                Shape::M16n8k256 => {
1083                    push_directive(tokens, "m16n8k256");
1084                }
1085                Shape::M8n8k128 => {
1086                    push_directive(tokens, "m8n8k128");
1087                }
1088            }
1089            push_directive(tokens, "row");
1090            push_directive(tokens, "col");
1091            push_directive(tokens, "s32");
1092            push_directive(tokens, "b1");
1093            push_directive(tokens, "b1");
1094            push_directive(tokens, "s32");
1095            match &self.bitop {
1096                Bitop::Xor => {
1097                    push_directive(tokens, "xor");
1098                }
1099                Bitop::And => {
1100                    push_directive(tokens, "and");
1101                }
1102            }
1103            push_directive(tokens, "popc");
1104            if spaced {
1105                tokens.push(PtxToken::Space);
1106            }
1107            self.d.unparse_tokens_mode(tokens, spaced);
1108            tokens.push(PtxToken::Comma);
1109            if spaced {
1110                tokens.push(PtxToken::Space);
1111            }
1112            self.a.unparse_tokens_mode(tokens, spaced);
1113            tokens.push(PtxToken::Comma);
1114            if spaced {
1115                tokens.push(PtxToken::Space);
1116            }
1117            self.b.unparse_tokens_mode(tokens, spaced);
1118            tokens.push(PtxToken::Comma);
1119            if spaced {
1120                tokens.push(PtxToken::Space);
1121            }
1122            self.c.unparse_tokens_mode(tokens, spaced);
1123            tokens.push(PtxToken::Semicolon);
1124            if spaced {
1125                tokens.push(PtxToken::Newline);
1126            }
1127        }
1128    }
1129}
ptx_parser/unparser/instruction/mma.rs

ptx_parser/unparser/instruction/
mma.rs
