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

ptx_parser/unparser/instruction/
mma.rs
