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            push_opcode(tokens, "mma");
84            push_directive(tokens, "sync");
85            push_directive(tokens, "aligned");
86            push_directive(tokens, "m8n8k4");
87            match &self.alayout {
88                Alayout::Row => {
89                    push_directive(tokens, "row");
90                }
91                Alayout::Col => {
92                    push_directive(tokens, "col");
93                }
94            }
95            match &self.blayout {
96                Blayout::Row => {
97                    push_directive(tokens, "row");
98                }
99                Blayout::Col => {
100                    push_directive(tokens, "col");
101                }
102            }
103            match &self.dtype {
104                Dtype::F16 => {
105                    push_directive(tokens, "f16");
106                }
107                Dtype::F32 => {
108                    push_directive(tokens, "f32");
109                }
110            }
111            push_directive(tokens, "f16");
112            push_directive(tokens, "f16");
113            match &self.ctype {
114                Ctype::F16 => {
115                    push_directive(tokens, "f16");
116                }
117                Ctype::F32 => {
118                    push_directive(tokens, "f32");
119                }
120            }
121            self.d.unparse_tokens(tokens);
122            tokens.push(PtxToken::Comma);
123            self.a.unparse_tokens(tokens);
124            tokens.push(PtxToken::Comma);
125            self.b.unparse_tokens(tokens);
126            tokens.push(PtxToken::Comma);
127            self.c.unparse_tokens(tokens);
128            tokens.push(PtxToken::Semicolon);
129        }
130    }
131
132    impl PtxUnparser for MmaSyncAlignedM16n8k8RowColDtypeF16F16Ctype {
133        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
134            push_opcode(tokens, "mma");
135            push_directive(tokens, "sync");
136            push_directive(tokens, "aligned");
137            push_directive(tokens, "m16n8k8");
138            push_directive(tokens, "row");
139            push_directive(tokens, "col");
140            match &self.dtype {
141                Dtype::F16 => {
142                    push_directive(tokens, "f16");
143                }
144                Dtype::F32 => {
145                    push_directive(tokens, "f32");
146                }
147            }
148            push_directive(tokens, "f16");
149            push_directive(tokens, "f16");
150            match &self.ctype {
151                Ctype::F16 => {
152                    push_directive(tokens, "f16");
153                }
154                Ctype::F32 => {
155                    push_directive(tokens, "f32");
156                }
157            }
158            self.d.unparse_tokens(tokens);
159            tokens.push(PtxToken::Comma);
160            self.a.unparse_tokens(tokens);
161            tokens.push(PtxToken::Comma);
162            self.b.unparse_tokens(tokens);
163            tokens.push(PtxToken::Comma);
164            self.c.unparse_tokens(tokens);
165            tokens.push(PtxToken::Semicolon);
166        }
167    }
168
169    impl PtxUnparser for MmaSyncAlignedM16n8k16RowColDtypeF16F16Ctype {
170        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
171            push_opcode(tokens, "mma");
172            push_directive(tokens, "sync");
173            push_directive(tokens, "aligned");
174            push_directive(tokens, "m16n8k16");
175            push_directive(tokens, "row");
176            push_directive(tokens, "col");
177            match &self.dtype {
178                Dtype::F16 => {
179                    push_directive(tokens, "f16");
180                }
181                Dtype::F32 => {
182                    push_directive(tokens, "f32");
183                }
184            }
185            push_directive(tokens, "f16");
186            push_directive(tokens, "f16");
187            match &self.ctype {
188                Ctype::F16 => {
189                    push_directive(tokens, "f16");
190                }
191                Ctype::F32 => {
192                    push_directive(tokens, "f32");
193                }
194            }
195            self.d.unparse_tokens(tokens);
196            tokens.push(PtxToken::Comma);
197            self.a.unparse_tokens(tokens);
198            tokens.push(PtxToken::Comma);
199            self.b.unparse_tokens(tokens);
200            tokens.push(PtxToken::Comma);
201            self.c.unparse_tokens(tokens);
202            tokens.push(PtxToken::Semicolon);
203        }
204    }
205}
206
207pub mod section_1 {
208    use super::*;
209    use crate::r#type::instruction::mma::section_1::*;
210
211    impl PtxUnparser for MmaSyncAlignedM16n8k4RowColF32Tf32Tf32F32 {
212        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
213            push_opcode(tokens, "mma");
214            push_directive(tokens, "sync");
215            push_directive(tokens, "aligned");
216            push_directive(tokens, "m16n8k4");
217            push_directive(tokens, "row");
218            push_directive(tokens, "col");
219            push_directive(tokens, "f32");
220            push_directive(tokens, "tf32");
221            push_directive(tokens, "tf32");
222            push_directive(tokens, "f32");
223            self.d.unparse_tokens(tokens);
224            tokens.push(PtxToken::Comma);
225            self.a.unparse_tokens(tokens);
226            tokens.push(PtxToken::Comma);
227            self.b.unparse_tokens(tokens);
228            tokens.push(PtxToken::Comma);
229            self.c.unparse_tokens(tokens);
230            tokens.push(PtxToken::Semicolon);
231        }
232    }
233
234    impl PtxUnparser for MmaSyncAlignedM16n8k8RowColF32AtypeBtypeF32 {
235        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
236            push_opcode(tokens, "mma");
237            push_directive(tokens, "sync");
238            push_directive(tokens, "aligned");
239            push_directive(tokens, "m16n8k8");
240            push_directive(tokens, "row");
241            push_directive(tokens, "col");
242            push_directive(tokens, "f32");
243            match &self.atype {
244                Atype::Bf16 => {
245                    push_directive(tokens, "bf16");
246                }
247                Atype::Tf32 => {
248                    push_directive(tokens, "tf32");
249                }
250            }
251            match &self.btype {
252                Btype::Bf16 => {
253                    push_directive(tokens, "bf16");
254                }
255                Btype::Tf32 => {
256                    push_directive(tokens, "tf32");
257                }
258            }
259            push_directive(tokens, "f32");
260            self.d.unparse_tokens(tokens);
261            tokens.push(PtxToken::Comma);
262            self.a.unparse_tokens(tokens);
263            tokens.push(PtxToken::Comma);
264            self.b.unparse_tokens(tokens);
265            tokens.push(PtxToken::Comma);
266            self.c.unparse_tokens(tokens);
267            tokens.push(PtxToken::Semicolon);
268        }
269    }
270
271    impl PtxUnparser for MmaSyncAlignedM16n8k16RowColF32Bf16Bf16F32 {
272        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
273            push_opcode(tokens, "mma");
274            push_directive(tokens, "sync");
275            push_directive(tokens, "aligned");
276            push_directive(tokens, "m16n8k16");
277            push_directive(tokens, "row");
278            push_directive(tokens, "col");
279            push_directive(tokens, "f32");
280            push_directive(tokens, "bf16");
281            push_directive(tokens, "bf16");
282            push_directive(tokens, "f32");
283            self.d.unparse_tokens(tokens);
284            tokens.push(PtxToken::Comma);
285            self.a.unparse_tokens(tokens);
286            tokens.push(PtxToken::Comma);
287            self.b.unparse_tokens(tokens);
288            tokens.push(PtxToken::Comma);
289            self.c.unparse_tokens(tokens);
290            tokens.push(PtxToken::Semicolon);
291        }
292    }
293
294    impl PtxUnparser for MmaSyncAlignedShapeRowColDtypeF8typeF8typeCtype {
295        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
296            push_opcode(tokens, "mma");
297            push_directive(tokens, "sync");
298            push_directive(tokens, "aligned");
299            match &self.shape {
300                Shape::M16n8k16 => {
301                    push_directive(tokens, "m16n8k16");
302                }
303                Shape::M16n8k32 => {
304                    push_directive(tokens, "m16n8k32");
305                }
306            }
307            push_directive(tokens, "row");
308            push_directive(tokens, "col");
309            match &self.dtype {
310                Dtype::F16 => {
311                    push_directive(tokens, "f16");
312                }
313                Dtype::F32 => {
314                    push_directive(tokens, "f32");
315                }
316            }
317            match &self.f8type {
318                F8type::E4m3 => {
319                    push_directive(tokens, "e4m3");
320                }
321                F8type::E5m2 => {
322                    push_directive(tokens, "e5m2");
323                }
324            }
325            match &self.f8type1 {
326                F8type::E4m3 => {
327                    push_directive(tokens, "e4m3");
328                }
329                F8type::E5m2 => {
330                    push_directive(tokens, "e5m2");
331                }
332            }
333            match &self.ctype {
334                Ctype::F16 => {
335                    push_directive(tokens, "f16");
336                }
337                Ctype::F32 => {
338                    push_directive(tokens, "f32");
339                }
340            }
341            self.d.unparse_tokens(tokens);
342            tokens.push(PtxToken::Comma);
343            self.a.unparse_tokens(tokens);
344            tokens.push(PtxToken::Comma);
345            self.b.unparse_tokens(tokens);
346            tokens.push(PtxToken::Comma);
347            self.c.unparse_tokens(tokens);
348            tokens.push(PtxToken::Semicolon);
349        }
350    }
351
352    impl PtxUnparser for MmaSyncAlignedM16n8k32RowColKindDtypeF8f6f4typeF8f6f4typeCtype {
353        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
354            push_opcode(tokens, "mma");
355            push_directive(tokens, "sync");
356            push_directive(tokens, "aligned");
357            push_directive(tokens, "m16n8k32");
358            push_directive(tokens, "row");
359            push_directive(tokens, "col");
360            match &self.kind {
361                Kind::KindF8f6f4 => {
362                    push_directive(tokens, "kind::f8f6f4");
363                }
364            }
365            match &self.dtype {
366                Dtype::F16 => {
367                    push_directive(tokens, "f16");
368                }
369                Dtype::F32 => {
370                    push_directive(tokens, "f32");
371                }
372            }
373            match &self.f8f6f4type {
374                F8f6f4type::E4m3 => {
375                    push_directive(tokens, "e4m3");
376                }
377                F8f6f4type::E5m2 => {
378                    push_directive(tokens, "e5m2");
379                }
380                F8f6f4type::E3m2 => {
381                    push_directive(tokens, "e3m2");
382                }
383                F8f6f4type::E2m3 => {
384                    push_directive(tokens, "e2m3");
385                }
386                F8f6f4type::E2m1 => {
387                    push_directive(tokens, "e2m1");
388                }
389            }
390            match &self.f8f6f4type1 {
391                F8f6f4type::E4m3 => {
392                    push_directive(tokens, "e4m3");
393                }
394                F8f6f4type::E5m2 => {
395                    push_directive(tokens, "e5m2");
396                }
397                F8f6f4type::E3m2 => {
398                    push_directive(tokens, "e3m2");
399                }
400                F8f6f4type::E2m3 => {
401                    push_directive(tokens, "e2m3");
402                }
403                F8f6f4type::E2m1 => {
404                    push_directive(tokens, "e2m1");
405                }
406            }
407            match &self.ctype {
408                Ctype::F16 => {
409                    push_directive(tokens, "f16");
410                }
411                Ctype::F32 => {
412                    push_directive(tokens, "f32");
413                }
414            }
415            self.d.unparse_tokens(tokens);
416            tokens.push(PtxToken::Comma);
417            self.a.unparse_tokens(tokens);
418            tokens.push(PtxToken::Comma);
419            self.b.unparse_tokens(tokens);
420            tokens.push(PtxToken::Comma);
421            self.c.unparse_tokens(tokens);
422            tokens.push(PtxToken::Semicolon);
423        }
424    }
425}
426
427pub mod section_2 {
428    use super::*;
429    use crate::r#type::instruction::mma::section_2::*;
430
431    impl PtxUnparser for MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype {
432        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
433            push_opcode(tokens, "mma");
434            push_directive(tokens, "sync");
435            push_directive(tokens, "aligned");
436            push_directive(tokens, "m16n8k64");
437            push_directive(tokens, "row");
438            push_directive(tokens, "col");
439            match &self.kind {
440                Kind::KindMxf4 => {
441                    push_directive(tokens, "kind::mxf4");
442                }
443            }
444            push_directive(tokens, "block_scale");
445            if let Some(scale_vec_size_0) = self.scale_vec_size.as_ref() {
446                match scale_vec_size_0 {
447                    ScaleVecSize::ScaleVec2x => {
448                        push_directive(tokens, "scale_vec::2X");
449                    }
450                }
451            }
452            push_directive(tokens, "f32");
453            push_directive(tokens, "e2m1");
454            push_directive(tokens, "e2m1");
455            push_directive(tokens, "f32");
456            match &self.stype {
457                Stype::Ue8m0 => {
458                    push_directive(tokens, "ue8m0");
459                }
460            }
461            self.d.unparse_tokens(tokens);
462            tokens.push(PtxToken::Comma);
463            self.a.unparse_tokens(tokens);
464            tokens.push(PtxToken::Comma);
465            self.b.unparse_tokens(tokens);
466            tokens.push(PtxToken::Comma);
467            self.c.unparse_tokens(tokens);
468            tokens.push(PtxToken::Comma);
469            self.scale_a_data.unparse_tokens(tokens);
470            tokens.push(PtxToken::Comma);
471            self.byte_id_a.unparse_tokens(tokens);
472            tokens.push(PtxToken::Comma);
473            self.scale_b_data.unparse_tokens(tokens);
474            tokens.push(PtxToken::Comma);
475            self.byte_id_b.unparse_tokens(tokens);
476            tokens.push(PtxToken::Semicolon);
477        }
478    }
479}
480
481pub mod section_3 {
482    use super::*;
483    use crate::r#type::instruction::mma::section_3::*;
484
485    impl PtxUnparser for MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype1 {
486        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
487            push_opcode(tokens, "mma");
488            push_directive(tokens, "sync");
489            push_directive(tokens, "aligned");
490            push_directive(tokens, "m16n8k64");
491            push_directive(tokens, "row");
492            push_directive(tokens, "col");
493            match &self.kind {
494                Kind::KindMxf4nvf4 => {
495                    push_directive(tokens, "kind::mxf4nvf4");
496                }
497            }
498            push_directive(tokens, "block_scale");
499            match &self.scale_vec_size {
500                ScaleVecSize::ScaleVec2x => {
501                    push_directive(tokens, "scale_vec::2X");
502                }
503                ScaleVecSize::ScaleVec4x => {
504                    push_directive(tokens, "scale_vec::4X");
505                }
506            }
507            push_directive(tokens, "f32");
508            push_directive(tokens, "e2m1");
509            push_directive(tokens, "e2m1");
510            push_directive(tokens, "f32");
511            match &self.stype {
512                Stype::Ue8m0 => {
513                    push_directive(tokens, "ue8m0");
514                }
515                Stype::Ue4m3 => {
516                    push_directive(tokens, "ue4m3");
517                }
518            }
519            self.d.unparse_tokens(tokens);
520            tokens.push(PtxToken::Comma);
521            self.a.unparse_tokens(tokens);
522            tokens.push(PtxToken::Comma);
523            self.b.unparse_tokens(tokens);
524            tokens.push(PtxToken::Comma);
525            self.c.unparse_tokens(tokens);
526            tokens.push(PtxToken::Comma);
527            self.scale_a_data.unparse_tokens(tokens);
528            tokens.push(PtxToken::Comma);
529            self.byte_id_a.unparse_tokens(tokens);
530            tokens.push(PtxToken::Comma);
531            self.scale_b_data.unparse_tokens(tokens);
532            tokens.push(PtxToken::Comma);
533            self.byte_id_b.unparse_tokens(tokens);
534            tokens.push(PtxToken::Semicolon);
535        }
536    }
537}
538
539pub mod section_4 {
540    use super::*;
541    use crate::r#type::instruction::mma::section_4::*;
542
543    impl PtxUnparser
544        for MmaSyncAlignedM16n8k32RowColKindBlockScaleScaleVecSizeF32F8f6f4typeF8f6f4typeF32Stype
545    {
546        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
547            push_opcode(tokens, "mma");
548            push_directive(tokens, "sync");
549            push_directive(tokens, "aligned");
550            push_directive(tokens, "m16n8k32");
551            push_directive(tokens, "row");
552            push_directive(tokens, "col");
553            match &self.kind {
554                Kind::KindMxf8f6f4 => {
555                    push_directive(tokens, "kind::mxf8f6f4");
556                }
557            }
558            push_directive(tokens, "block_scale");
559            if let Some(scale_vec_size_1) = self.scale_vec_size.as_ref() {
560                match scale_vec_size_1 {
561                    ScaleVecSize::ScaleVec1x => {
562                        push_directive(tokens, "scale_vec::1X");
563                    }
564                }
565            }
566            push_directive(tokens, "f32");
567            match &self.f8f6f4type {
568                F8f6f4type::E4m3 => {
569                    push_directive(tokens, "e4m3");
570                }
571                F8f6f4type::E5m2 => {
572                    push_directive(tokens, "e5m2");
573                }
574                F8f6f4type::E3m2 => {
575                    push_directive(tokens, "e3m2");
576                }
577                F8f6f4type::E2m3 => {
578                    push_directive(tokens, "e2m3");
579                }
580                F8f6f4type::E2m1 => {
581                    push_directive(tokens, "e2m1");
582                }
583            }
584            match &self.f8f6f4type1 {
585                F8f6f4type::E4m3 => {
586                    push_directive(tokens, "e4m3");
587                }
588                F8f6f4type::E5m2 => {
589                    push_directive(tokens, "e5m2");
590                }
591                F8f6f4type::E3m2 => {
592                    push_directive(tokens, "e3m2");
593                }
594                F8f6f4type::E2m3 => {
595                    push_directive(tokens, "e2m3");
596                }
597                F8f6f4type::E2m1 => {
598                    push_directive(tokens, "e2m1");
599                }
600            }
601            push_directive(tokens, "f32");
602            match &self.stype {
603                Stype::Ue8m0 => {
604                    push_directive(tokens, "ue8m0");
605                }
606            }
607            self.d.unparse_tokens(tokens);
608            tokens.push(PtxToken::Comma);
609            self.a.unparse_tokens(tokens);
610            tokens.push(PtxToken::Comma);
611            self.b.unparse_tokens(tokens);
612            tokens.push(PtxToken::Comma);
613            self.c.unparse_tokens(tokens);
614            tokens.push(PtxToken::Comma);
615            self.scale_a_data.unparse_tokens(tokens);
616            tokens.push(PtxToken::Comma);
617            self.byte_id_a.unparse_tokens(tokens);
618            tokens.push(PtxToken::Comma);
619            self.scale_b_data.unparse_tokens(tokens);
620            tokens.push(PtxToken::Comma);
621            self.byte_id_b.unparse_tokens(tokens);
622            tokens.push(PtxToken::Semicolon);
623        }
624    }
625}
626
627pub mod section_5 {
628    use super::*;
629    use crate::r#type::instruction::mma::section_5::*;
630
631    impl PtxUnparser for MmaSyncAlignedShapeRowColF64F64F64F64 {
632        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
633            push_opcode(tokens, "mma");
634            push_directive(tokens, "sync");
635            push_directive(tokens, "aligned");
636            match &self.shape {
637                Shape::M16n8k16 => {
638                    push_directive(tokens, "m16n8k16");
639                }
640                Shape::M16n8k4 => {
641                    push_directive(tokens, "m16n8k4");
642                }
643                Shape::M16n8k8 => {
644                    push_directive(tokens, "m16n8k8");
645                }
646                Shape::M8n84 => {
647                    push_directive(tokens, "m8n84");
648                }
649            }
650            push_directive(tokens, "row");
651            push_directive(tokens, "col");
652            push_directive(tokens, "f64");
653            push_directive(tokens, "f64");
654            push_directive(tokens, "f64");
655            push_directive(tokens, "f64");
656            self.d.unparse_tokens(tokens);
657            tokens.push(PtxToken::Comma);
658            self.a.unparse_tokens(tokens);
659            tokens.push(PtxToken::Comma);
660            self.b.unparse_tokens(tokens);
661            tokens.push(PtxToken::Comma);
662            self.c.unparse_tokens(tokens);
663            tokens.push(PtxToken::Semicolon);
664        }
665    }
666}
667
668pub mod section_6 {
669    use super::*;
670    use crate::r#type::instruction::mma::section_6::*;
671
672    impl PtxUnparser for MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS32 {
673        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
674            push_opcode(tokens, "mma");
675            push_directive(tokens, "sync");
676            push_directive(tokens, "aligned");
677            match &self.shape {
678                Shape::M16n8k16 => {
679                    push_directive(tokens, "m16n8k16");
680                }
681                Shape::M16n8k32 => {
682                    push_directive(tokens, "m16n8k32");
683                }
684                Shape::M8n8k16 => {
685                    push_directive(tokens, "m8n8k16");
686                }
687            }
688            push_directive(tokens, "row");
689            push_directive(tokens, "col");
690            if self.satfinite {
691                push_directive(tokens, "satfinite");
692            }
693            push_directive(tokens, "s32");
694            match &self.atype {
695                Atype::U8 => {
696                    push_directive(tokens, "u8");
697                }
698                Atype::S8 => {
699                    push_directive(tokens, "s8");
700                }
701            }
702            match &self.btype {
703                Btype::U8 => {
704                    push_directive(tokens, "u8");
705                }
706                Btype::S8 => {
707                    push_directive(tokens, "s8");
708                }
709            }
710            push_directive(tokens, "s32");
711            self.d.unparse_tokens(tokens);
712            tokens.push(PtxToken::Comma);
713            self.a.unparse_tokens(tokens);
714            tokens.push(PtxToken::Comma);
715            self.b.unparse_tokens(tokens);
716            tokens.push(PtxToken::Comma);
717            self.c.unparse_tokens(tokens);
718            tokens.push(PtxToken::Semicolon);
719        }
720    }
721}
722
723pub mod section_7 {
724    use super::*;
725    use crate::r#type::instruction::mma::section_7::*;
726
727    impl PtxUnparser for MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS321 {
728        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
729            push_opcode(tokens, "mma");
730            push_directive(tokens, "sync");
731            push_directive(tokens, "aligned");
732            match &self.shape {
733                Shape::M16n8k32 => {
734                    push_directive(tokens, "m16n8k32");
735                }
736                Shape::M16n8k64 => {
737                    push_directive(tokens, "m16n8k64");
738                }
739                Shape::M8n8k32 => {
740                    push_directive(tokens, "m8n8k32");
741                }
742            }
743            push_directive(tokens, "row");
744            push_directive(tokens, "col");
745            if self.satfinite {
746                push_directive(tokens, "satfinite");
747            }
748            push_directive(tokens, "s32");
749            match &self.atype {
750                Atype::U4 => {
751                    push_directive(tokens, "u4");
752                }
753                Atype::S4 => {
754                    push_directive(tokens, "s4");
755                }
756            }
757            match &self.btype {
758                Btype::U4 => {
759                    push_directive(tokens, "u4");
760                }
761                Btype::S4 => {
762                    push_directive(tokens, "s4");
763                }
764            }
765            push_directive(tokens, "s32");
766            self.d.unparse_tokens(tokens);
767            tokens.push(PtxToken::Comma);
768            self.a.unparse_tokens(tokens);
769            tokens.push(PtxToken::Comma);
770            self.b.unparse_tokens(tokens);
771            tokens.push(PtxToken::Comma);
772            self.c.unparse_tokens(tokens);
773            tokens.push(PtxToken::Semicolon);
774        }
775    }
776}
777
778pub mod section_8 {
779    use super::*;
780    use crate::r#type::instruction::mma::section_8::*;
781
782    impl PtxUnparser for MmaSyncAlignedShapeRowColS32B1B1S32BitopPopc {
783        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
784            push_opcode(tokens, "mma");
785            push_directive(tokens, "sync");
786            push_directive(tokens, "aligned");
787            match &self.shape {
788                Shape::M16n8k128 => {
789                    push_directive(tokens, "m16n8k128");
790                }
791                Shape::M16n8k256 => {
792                    push_directive(tokens, "m16n8k256");
793                }
794                Shape::M8n8k128 => {
795                    push_directive(tokens, "m8n8k128");
796                }
797            }
798            push_directive(tokens, "row");
799            push_directive(tokens, "col");
800            push_directive(tokens, "s32");
801            push_directive(tokens, "b1");
802            push_directive(tokens, "b1");
803            push_directive(tokens, "s32");
804            match &self.bitop {
805                Bitop::Xor => {
806                    push_directive(tokens, "xor");
807                }
808                Bitop::And => {
809                    push_directive(tokens, "and");
810                }
811            }
812            push_directive(tokens, "popc");
813            self.d.unparse_tokens(tokens);
814            tokens.push(PtxToken::Comma);
815            self.a.unparse_tokens(tokens);
816            tokens.push(PtxToken::Comma);
817            self.b.unparse_tokens(tokens);
818            tokens.push(PtxToken::Comma);
819            self.c.unparse_tokens(tokens);
820            tokens.push(PtxToken::Semicolon);
821        }
822    }
823}
ptx_parser/unparser/instruction/mma.rs

ptx_parser/unparser/instruction/
mma.rs
