ptx_parser/unparser/instruction/
wgmma_mma_async.rs

1//! Original PTX specification:
2//!
3//! // Half precision floating point type:
4//! wgmma.mma_async.sync.aligned.shape.dtype.f16.f16  d, a-desc, b-desc, scale-d, imm-scale-a, imm-scale-b, imm-trans-a, imm-trans-b;
5//! wgmma.mma_async.sync.aligned.shape.dtype.f16.f16  d, a, b-desc, scale-d, imm-scale-a, imm-scale-b, imm-trans-b;
6//! .shape   = {.m64n8k16, .m64n16k16, .m64n24k16, .m64n32k16,
7//! .m64n40k16, .m64n48k16, .m64n56k16, .m64n64k16,
8//! .m64n72k16, .m64n80k16, .m64n88k16, .m64n96k16,
9//! .m64n104k16, .m64n112k16, .m64n120k16, .m64n128k16,
10//! .m64n136k16, .m64n144k16, .m64n152k16, .m64n160k16,
11//! .m64n168k16, .m64n176k16, .m64n184k16, .m64n192k16,
12//! .m64n200k16, .m64n208k16, .m64n216k16, .m64n224k16,
13//! .m64n232k16, .m64n240k16, .m64n248k16, .m64n256k16};
14//! .dtype   = {.f16, .f32};
15//! ------------------------------------------------------------------
16//! // Alternate floating point type :
17//! // .bf16 floating point type:
18//! wgmma.mma_async.sync.aligned.shape.dtype.bf16.bf16  d, a-desc, b-desc, scale-d, imm-scale-a, imm-scale-b, imm-trans-a, imm-trans-b;
19//! wgmma.mma_async.sync.aligned.shape.dtype.bf16.bf16  d, a, b-desc, scale-d, imm-scale-a, imm-scale-b, imm-trans-b;
20//! .shape   = {.m64n8k16, .m64n16k16, .m64n24k16, .m64n32k16,
21//! .m64n40k16, .m64n48k16, .m64n56k16, .m64n64k16,
22//! .m64n72k16, .m64n80k16, .m64n88k16, .m64n96k16,
23//! .m64n104k16, .m64n112k16, .m64n120k16, .m64n128k16,
24//! .m64n136k16, .m64n144k16, .m64n152k16, .m64n160k16,
25//! .m64n168k16, .m64n176k16, .m64n184k16, .m64n192k16,
26//! .m64n200k16, .m64n208k16, .m64n216k16, .m64n224k16,
27//! .m64n232k16, .m64n240k16, .m64n248k16, .m64n256k16};
28//! .dtype  = {.f32};
29//! ------------------------------------------------------------------
30//! // .tf32 floating point type:
31//! wgmma.mma_async.sync.aligned.shape.dtype.tf32.tf32  d, a-desc, b-desc, scale-d, imm-scale-a, imm-scale-b;
32//! wgmma.mma_async.sync.aligned.shape.dtype.tf32.tf32  d, a, b-desc, scale-d, imm-scale-a, imm-scale-b;
33//! .shape   = {.m64n8k8, .m64n16k8, .m64n24k8, .m64n32k8,
34//! .m64n40k8, .m64n48k8, .m64n56k8, .m64n64k8,
35//! .m64n72k8, .m64n80k8, .m64n88k8, .m64n96k8,
36//! .m64n104k8, .m64n112k8, .m64n120k8, .m64n128k8,
37//! .m64n136k8, .m64n144k8, .m64n152k8, .m64n160k8,
38//! .m64n168k8, .m64n176k8, .m64n184k8, .m64n192k8,
39//! .m64n200k8, .m64n208k8, .m64n216k8, .m64n224k8,
40//! .m64n232k8, .m64n240k8, .m64n248k8, .m64n256k8};
41//! .dtype  = {.f32};
42//! ------------------------------------------------------------------
43//! // FP8 floating point type
44//! wgmma.mma_async.sync.aligned.shape.dtype.atype.btype  d, a-desc, b-desc, scale-d, imm-scale-a, imm-scale-b;
45//! wgmma.mma_async.sync.aligned.shape.dtype.atype.btype  d, a, b-desc, scale-d, imm-scale-a, imm-scale-b;
46//! .shape   = {.m64n8k32, .m64n16k32, .m64n24k32, .m64n32k32,
47//! .m64n40k32, .m64n48k32, .m64n56k32, .m64n64k32,
48//! .m64n72k32, .m64n80k32, .m64n88k32, .m64n96k32,
49//! .m64n104k32, .m64n112k32, .m64n120k32, .m64n128k32,
50//! .m64n136k32, .m64n144k32, .m64n152k32, .m64n160k32,
51//! .m64n168k32, .m64n176k32, .m64n184k32, .m64n192k32,
52//! .m64n200k32, .m64n208k32, .m64n216k32, .m64n224k32,
53//! .m64n232k32, .m64n240k32, .m64n248k32, .m64n256k32};
54//! .atype  = {.e4m3, .e5m2};
55//! .btype  = {.e4m3, .e5m2};
56//! .dtype  = {.f16, .f32};
57//! ------------------------------------------------------------------
58//! // Integer type:
59//! wgmma.mma_async.sync.aligned.shape{.satfinite}.s32.atype.btype  d, a-desc, b-desc, scale-d;
60//! wgmma.mma_async.sync.aligned.shape{.satfinite}.s32.atype.btype  d, a, b-desc, scale-d;
61//! .shape   = {.m64n8k32, .m64n16k32, .m64n24k32, .m64n32k32,
62//! .m64n48k32, .m64n64k32, .m64n80k32, .m64n96k32,
63//! .m64n112k32, .m64n128k32, .m64n144k32, .m64n160k32,
64//! .m64n176k32, .m64n192k32, .m64n208k32, .m64n224k32};
65//! .atype  = {.s8, .u8};
66//! .btype  = {.s8, .u8};
67//! ------------------------------------------------------------------
68//! // Single bit:
69//! wgmma.mma_async.sync.aligned.shape.s32.b1.b1.op.popc  d, a-desc, b-desc, scale-d;
70//! wgmma.mma_async.sync.aligned.shape.s32.b1.b1.op.popc  d, a, b-desc, scale-d;
71//! .shape   = {.m64n8k256, .m64n16k256, .m64n24k256, .m64n32k256,
72//! .m64n48k256, .m64n64k256, .m64n80k256, .m64n96k256,
73//! .m64n112k256, .m64n128k256, .m64n144k256, .m64n160k256,
74//! .m64n176k256, .m64n192k256, .m64n208k256, .m64n224k256,
75//! .m64n240k256, .m64n256k256};
76//! .op  = {.and};
77
78#![allow(unused)]
79
80use crate::lexer::PtxToken;
81use crate::unparser::{PtxUnparser, common::*};
82
83pub mod section_0 {
84    use super::*;
85    use crate::r#type::instruction::wgmma_mma_async::section_0::*;
86
87    impl PtxUnparser for WgmmaMmaAsyncSyncAlignedShapeDtypeF16F16 {
88        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
89            self.unparse_tokens_mode(tokens, false);
90        }
91        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
92            push_opcode(tokens, "wgmma");
93            push_directive(tokens, "mma_async");
94            push_directive(tokens, "sync");
95            push_directive(tokens, "aligned");
96            match &self.shape {
97                Shape::M64n104k16 => {
98                    push_directive(tokens, "m64n104k16");
99                }
100                Shape::M64n112k16 => {
101                    push_directive(tokens, "m64n112k16");
102                }
103                Shape::M64n120k16 => {
104                    push_directive(tokens, "m64n120k16");
105                }
106                Shape::M64n128k16 => {
107                    push_directive(tokens, "m64n128k16");
108                }
109                Shape::M64n136k16 => {
110                    push_directive(tokens, "m64n136k16");
111                }
112                Shape::M64n144k16 => {
113                    push_directive(tokens, "m64n144k16");
114                }
115                Shape::M64n152k16 => {
116                    push_directive(tokens, "m64n152k16");
117                }
118                Shape::M64n160k16 => {
119                    push_directive(tokens, "m64n160k16");
120                }
121                Shape::M64n168k16 => {
122                    push_directive(tokens, "m64n168k16");
123                }
124                Shape::M64n176k16 => {
125                    push_directive(tokens, "m64n176k16");
126                }
127                Shape::M64n184k16 => {
128                    push_directive(tokens, "m64n184k16");
129                }
130                Shape::M64n192k16 => {
131                    push_directive(tokens, "m64n192k16");
132                }
133                Shape::M64n200k16 => {
134                    push_directive(tokens, "m64n200k16");
135                }
136                Shape::M64n208k16 => {
137                    push_directive(tokens, "m64n208k16");
138                }
139                Shape::M64n216k16 => {
140                    push_directive(tokens, "m64n216k16");
141                }
142                Shape::M64n224k16 => {
143                    push_directive(tokens, "m64n224k16");
144                }
145                Shape::M64n232k16 => {
146                    push_directive(tokens, "m64n232k16");
147                }
148                Shape::M64n240k16 => {
149                    push_directive(tokens, "m64n240k16");
150                }
151                Shape::M64n248k16 => {
152                    push_directive(tokens, "m64n248k16");
153                }
154                Shape::M64n256k16 => {
155                    push_directive(tokens, "m64n256k16");
156                }
157                Shape::M64n16k16 => {
158                    push_directive(tokens, "m64n16k16");
159                }
160                Shape::M64n24k16 => {
161                    push_directive(tokens, "m64n24k16");
162                }
163                Shape::M64n32k16 => {
164                    push_directive(tokens, "m64n32k16");
165                }
166                Shape::M64n40k16 => {
167                    push_directive(tokens, "m64n40k16");
168                }
169                Shape::M64n48k16 => {
170                    push_directive(tokens, "m64n48k16");
171                }
172                Shape::M64n56k16 => {
173                    push_directive(tokens, "m64n56k16");
174                }
175                Shape::M64n64k16 => {
176                    push_directive(tokens, "m64n64k16");
177                }
178                Shape::M64n72k16 => {
179                    push_directive(tokens, "m64n72k16");
180                }
181                Shape::M64n80k16 => {
182                    push_directive(tokens, "m64n80k16");
183                }
184                Shape::M64n88k16 => {
185                    push_directive(tokens, "m64n88k16");
186                }
187                Shape::M64n96k16 => {
188                    push_directive(tokens, "m64n96k16");
189                }
190                Shape::M64n8k16 => {
191                    push_directive(tokens, "m64n8k16");
192                }
193            }
194            match &self.dtype {
195                Dtype::F16 => {
196                    push_directive(tokens, "f16");
197                }
198                Dtype::F32 => {
199                    push_directive(tokens, "f32");
200                }
201            }
202            push_directive(tokens, "f16");
203            push_directive(tokens, "f16");
204            if spaced {
205                tokens.push(PtxToken::Space);
206            }
207            self.d.unparse_tokens_mode(tokens, spaced);
208            tokens.push(PtxToken::Comma);
209            if spaced {
210                tokens.push(PtxToken::Space);
211            }
212            self.a_desc.unparse_tokens_mode(tokens, spaced);
213            tokens.push(PtxToken::Comma);
214            if spaced {
215                tokens.push(PtxToken::Space);
216            }
217            self.b_desc.unparse_tokens_mode(tokens, spaced);
218            tokens.push(PtxToken::Comma);
219            if spaced {
220                tokens.push(PtxToken::Space);
221            }
222            self.scale_d.unparse_tokens_mode(tokens, spaced);
223            tokens.push(PtxToken::Comma);
224            if spaced {
225                tokens.push(PtxToken::Space);
226            }
227            self.imm_scale_a.unparse_tokens_mode(tokens, spaced);
228            tokens.push(PtxToken::Comma);
229            if spaced {
230                tokens.push(PtxToken::Space);
231            }
232            self.imm_scale_b.unparse_tokens_mode(tokens, spaced);
233            tokens.push(PtxToken::Comma);
234            if spaced {
235                tokens.push(PtxToken::Space);
236            }
237            self.imm_trans_a.unparse_tokens_mode(tokens, spaced);
238            tokens.push(PtxToken::Comma);
239            if spaced {
240                tokens.push(PtxToken::Space);
241            }
242            self.imm_trans_b.unparse_tokens_mode(tokens, spaced);
243            tokens.push(PtxToken::Semicolon);
244            if spaced {
245                tokens.push(PtxToken::Newline);
246            }
247        }
248    }
249
250    impl PtxUnparser for WgmmaMmaAsyncSyncAlignedShapeDtypeF16F161 {
251        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
252            self.unparse_tokens_mode(tokens, false);
253        }
254        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
255            push_opcode(tokens, "wgmma");
256            push_directive(tokens, "mma_async");
257            push_directive(tokens, "sync");
258            push_directive(tokens, "aligned");
259            match &self.shape {
260                Shape::M64n104k16 => {
261                    push_directive(tokens, "m64n104k16");
262                }
263                Shape::M64n112k16 => {
264                    push_directive(tokens, "m64n112k16");
265                }
266                Shape::M64n120k16 => {
267                    push_directive(tokens, "m64n120k16");
268                }
269                Shape::M64n128k16 => {
270                    push_directive(tokens, "m64n128k16");
271                }
272                Shape::M64n136k16 => {
273                    push_directive(tokens, "m64n136k16");
274                }
275                Shape::M64n144k16 => {
276                    push_directive(tokens, "m64n144k16");
277                }
278                Shape::M64n152k16 => {
279                    push_directive(tokens, "m64n152k16");
280                }
281                Shape::M64n160k16 => {
282                    push_directive(tokens, "m64n160k16");
283                }
284                Shape::M64n168k16 => {
285                    push_directive(tokens, "m64n168k16");
286                }
287                Shape::M64n176k16 => {
288                    push_directive(tokens, "m64n176k16");
289                }
290                Shape::M64n184k16 => {
291                    push_directive(tokens, "m64n184k16");
292                }
293                Shape::M64n192k16 => {
294                    push_directive(tokens, "m64n192k16");
295                }
296                Shape::M64n200k16 => {
297                    push_directive(tokens, "m64n200k16");
298                }
299                Shape::M64n208k16 => {
300                    push_directive(tokens, "m64n208k16");
301                }
302                Shape::M64n216k16 => {
303                    push_directive(tokens, "m64n216k16");
304                }
305                Shape::M64n224k16 => {
306                    push_directive(tokens, "m64n224k16");
307                }
308                Shape::M64n232k16 => {
309                    push_directive(tokens, "m64n232k16");
310                }
311                Shape::M64n240k16 => {
312                    push_directive(tokens, "m64n240k16");
313                }
314                Shape::M64n248k16 => {
315                    push_directive(tokens, "m64n248k16");
316                }
317                Shape::M64n256k16 => {
318                    push_directive(tokens, "m64n256k16");
319                }
320                Shape::M64n16k16 => {
321                    push_directive(tokens, "m64n16k16");
322                }
323                Shape::M64n24k16 => {
324                    push_directive(tokens, "m64n24k16");
325                }
326                Shape::M64n32k16 => {
327                    push_directive(tokens, "m64n32k16");
328                }
329                Shape::M64n40k16 => {
330                    push_directive(tokens, "m64n40k16");
331                }
332                Shape::M64n48k16 => {
333                    push_directive(tokens, "m64n48k16");
334                }
335                Shape::M64n56k16 => {
336                    push_directive(tokens, "m64n56k16");
337                }
338                Shape::M64n64k16 => {
339                    push_directive(tokens, "m64n64k16");
340                }
341                Shape::M64n72k16 => {
342                    push_directive(tokens, "m64n72k16");
343                }
344                Shape::M64n80k16 => {
345                    push_directive(tokens, "m64n80k16");
346                }
347                Shape::M64n88k16 => {
348                    push_directive(tokens, "m64n88k16");
349                }
350                Shape::M64n96k16 => {
351                    push_directive(tokens, "m64n96k16");
352                }
353                Shape::M64n8k16 => {
354                    push_directive(tokens, "m64n8k16");
355                }
356            }
357            match &self.dtype {
358                Dtype::F16 => {
359                    push_directive(tokens, "f16");
360                }
361                Dtype::F32 => {
362                    push_directive(tokens, "f32");
363                }
364            }
365            push_directive(tokens, "f16");
366            push_directive(tokens, "f16");
367            if spaced {
368                tokens.push(PtxToken::Space);
369            }
370            self.d.unparse_tokens_mode(tokens, spaced);
371            tokens.push(PtxToken::Comma);
372            if spaced {
373                tokens.push(PtxToken::Space);
374            }
375            self.a.unparse_tokens_mode(tokens, spaced);
376            tokens.push(PtxToken::Comma);
377            if spaced {
378                tokens.push(PtxToken::Space);
379            }
380            self.b_desc.unparse_tokens_mode(tokens, spaced);
381            tokens.push(PtxToken::Comma);
382            if spaced {
383                tokens.push(PtxToken::Space);
384            }
385            self.scale_d.unparse_tokens_mode(tokens, spaced);
386            tokens.push(PtxToken::Comma);
387            if spaced {
388                tokens.push(PtxToken::Space);
389            }
390            self.imm_scale_a.unparse_tokens_mode(tokens, spaced);
391            tokens.push(PtxToken::Comma);
392            if spaced {
393                tokens.push(PtxToken::Space);
394            }
395            self.imm_scale_b.unparse_tokens_mode(tokens, spaced);
396            tokens.push(PtxToken::Comma);
397            if spaced {
398                tokens.push(PtxToken::Space);
399            }
400            self.imm_trans_b.unparse_tokens_mode(tokens, spaced);
401            tokens.push(PtxToken::Semicolon);
402            if spaced {
403                tokens.push(PtxToken::Newline);
404            }
405        }
406    }
407}
408
409pub mod section_1 {
410    use super::*;
411    use crate::r#type::instruction::wgmma_mma_async::section_1::*;
412
413    impl PtxUnparser for WgmmaMmaAsyncSyncAlignedShapeDtypeBf16Bf16 {
414        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
415            self.unparse_tokens_mode(tokens, false);
416        }
417        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
418            push_opcode(tokens, "wgmma");
419            push_directive(tokens, "mma_async");
420            push_directive(tokens, "sync");
421            push_directive(tokens, "aligned");
422            match &self.shape {
423                Shape::M64n104k16 => {
424                    push_directive(tokens, "m64n104k16");
425                }
426                Shape::M64n112k16 => {
427                    push_directive(tokens, "m64n112k16");
428                }
429                Shape::M64n120k16 => {
430                    push_directive(tokens, "m64n120k16");
431                }
432                Shape::M64n128k16 => {
433                    push_directive(tokens, "m64n128k16");
434                }
435                Shape::M64n136k16 => {
436                    push_directive(tokens, "m64n136k16");
437                }
438                Shape::M64n144k16 => {
439                    push_directive(tokens, "m64n144k16");
440                }
441                Shape::M64n152k16 => {
442                    push_directive(tokens, "m64n152k16");
443                }
444                Shape::M64n160k16 => {
445                    push_directive(tokens, "m64n160k16");
446                }
447                Shape::M64n168k16 => {
448                    push_directive(tokens, "m64n168k16");
449                }
450                Shape::M64n176k16 => {
451                    push_directive(tokens, "m64n176k16");
452                }
453                Shape::M64n184k16 => {
454                    push_directive(tokens, "m64n184k16");
455                }
456                Shape::M64n192k16 => {
457                    push_directive(tokens, "m64n192k16");
458                }
459                Shape::M64n200k16 => {
460                    push_directive(tokens, "m64n200k16");
461                }
462                Shape::M64n208k16 => {
463                    push_directive(tokens, "m64n208k16");
464                }
465                Shape::M64n216k16 => {
466                    push_directive(tokens, "m64n216k16");
467                }
468                Shape::M64n224k16 => {
469                    push_directive(tokens, "m64n224k16");
470                }
471                Shape::M64n232k16 => {
472                    push_directive(tokens, "m64n232k16");
473                }
474                Shape::M64n240k16 => {
475                    push_directive(tokens, "m64n240k16");
476                }
477                Shape::M64n248k16 => {
478                    push_directive(tokens, "m64n248k16");
479                }
480                Shape::M64n256k16 => {
481                    push_directive(tokens, "m64n256k16");
482                }
483                Shape::M64n16k16 => {
484                    push_directive(tokens, "m64n16k16");
485                }
486                Shape::M64n24k16 => {
487                    push_directive(tokens, "m64n24k16");
488                }
489                Shape::M64n32k16 => {
490                    push_directive(tokens, "m64n32k16");
491                }
492                Shape::M64n40k16 => {
493                    push_directive(tokens, "m64n40k16");
494                }
495                Shape::M64n48k16 => {
496                    push_directive(tokens, "m64n48k16");
497                }
498                Shape::M64n56k16 => {
499                    push_directive(tokens, "m64n56k16");
500                }
501                Shape::M64n64k16 => {
502                    push_directive(tokens, "m64n64k16");
503                }
504                Shape::M64n72k16 => {
505                    push_directive(tokens, "m64n72k16");
506                }
507                Shape::M64n80k16 => {
508                    push_directive(tokens, "m64n80k16");
509                }
510                Shape::M64n88k16 => {
511                    push_directive(tokens, "m64n88k16");
512                }
513                Shape::M64n96k16 => {
514                    push_directive(tokens, "m64n96k16");
515                }
516                Shape::M64n8k16 => {
517                    push_directive(tokens, "m64n8k16");
518                }
519            }
520            match &self.dtype {
521                Dtype::F32 => {
522                    push_directive(tokens, "f32");
523                }
524            }
525            push_directive(tokens, "bf16");
526            push_directive(tokens, "bf16");
527            if spaced {
528                tokens.push(PtxToken::Space);
529            }
530            self.d.unparse_tokens_mode(tokens, spaced);
531            tokens.push(PtxToken::Comma);
532            if spaced {
533                tokens.push(PtxToken::Space);
534            }
535            self.a_desc.unparse_tokens_mode(tokens, spaced);
536            tokens.push(PtxToken::Comma);
537            if spaced {
538                tokens.push(PtxToken::Space);
539            }
540            self.b_desc.unparse_tokens_mode(tokens, spaced);
541            tokens.push(PtxToken::Comma);
542            if spaced {
543                tokens.push(PtxToken::Space);
544            }
545            self.scale_d.unparse_tokens_mode(tokens, spaced);
546            tokens.push(PtxToken::Comma);
547            if spaced {
548                tokens.push(PtxToken::Space);
549            }
550            self.imm_scale_a.unparse_tokens_mode(tokens, spaced);
551            tokens.push(PtxToken::Comma);
552            if spaced {
553                tokens.push(PtxToken::Space);
554            }
555            self.imm_scale_b.unparse_tokens_mode(tokens, spaced);
556            tokens.push(PtxToken::Comma);
557            if spaced {
558                tokens.push(PtxToken::Space);
559            }
560            self.imm_trans_a.unparse_tokens_mode(tokens, spaced);
561            tokens.push(PtxToken::Comma);
562            if spaced {
563                tokens.push(PtxToken::Space);
564            }
565            self.imm_trans_b.unparse_tokens_mode(tokens, spaced);
566            tokens.push(PtxToken::Semicolon);
567            if spaced {
568                tokens.push(PtxToken::Newline);
569            }
570        }
571    }
572
573    impl PtxUnparser for WgmmaMmaAsyncSyncAlignedShapeDtypeBf16Bf161 {
574        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
575            self.unparse_tokens_mode(tokens, false);
576        }
577        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
578            push_opcode(tokens, "wgmma");
579            push_directive(tokens, "mma_async");
580            push_directive(tokens, "sync");
581            push_directive(tokens, "aligned");
582            match &self.shape {
583                Shape::M64n104k16 => {
584                    push_directive(tokens, "m64n104k16");
585                }
586                Shape::M64n112k16 => {
587                    push_directive(tokens, "m64n112k16");
588                }
589                Shape::M64n120k16 => {
590                    push_directive(tokens, "m64n120k16");
591                }
592                Shape::M64n128k16 => {
593                    push_directive(tokens, "m64n128k16");
594                }
595                Shape::M64n136k16 => {
596                    push_directive(tokens, "m64n136k16");
597                }
598                Shape::M64n144k16 => {
599                    push_directive(tokens, "m64n144k16");
600                }
601                Shape::M64n152k16 => {
602                    push_directive(tokens, "m64n152k16");
603                }
604                Shape::M64n160k16 => {
605                    push_directive(tokens, "m64n160k16");
606                }
607                Shape::M64n168k16 => {
608                    push_directive(tokens, "m64n168k16");
609                }
610                Shape::M64n176k16 => {
611                    push_directive(tokens, "m64n176k16");
612                }
613                Shape::M64n184k16 => {
614                    push_directive(tokens, "m64n184k16");
615                }
616                Shape::M64n192k16 => {
617                    push_directive(tokens, "m64n192k16");
618                }
619                Shape::M64n200k16 => {
620                    push_directive(tokens, "m64n200k16");
621                }
622                Shape::M64n208k16 => {
623                    push_directive(tokens, "m64n208k16");
624                }
625                Shape::M64n216k16 => {
626                    push_directive(tokens, "m64n216k16");
627                }
628                Shape::M64n224k16 => {
629                    push_directive(tokens, "m64n224k16");
630                }
631                Shape::M64n232k16 => {
632                    push_directive(tokens, "m64n232k16");
633                }
634                Shape::M64n240k16 => {
635                    push_directive(tokens, "m64n240k16");
636                }
637                Shape::M64n248k16 => {
638                    push_directive(tokens, "m64n248k16");
639                }
640                Shape::M64n256k16 => {
641                    push_directive(tokens, "m64n256k16");
642                }
643                Shape::M64n16k16 => {
644                    push_directive(tokens, "m64n16k16");
645                }
646                Shape::M64n24k16 => {
647                    push_directive(tokens, "m64n24k16");
648                }
649                Shape::M64n32k16 => {
650                    push_directive(tokens, "m64n32k16");
651                }
652                Shape::M64n40k16 => {
653                    push_directive(tokens, "m64n40k16");
654                }
655                Shape::M64n48k16 => {
656                    push_directive(tokens, "m64n48k16");
657                }
658                Shape::M64n56k16 => {
659                    push_directive(tokens, "m64n56k16");
660                }
661                Shape::M64n64k16 => {
662                    push_directive(tokens, "m64n64k16");
663                }
664                Shape::M64n72k16 => {
665                    push_directive(tokens, "m64n72k16");
666                }
667                Shape::M64n80k16 => {
668                    push_directive(tokens, "m64n80k16");
669                }
670                Shape::M64n88k16 => {
671                    push_directive(tokens, "m64n88k16");
672                }
673                Shape::M64n96k16 => {
674                    push_directive(tokens, "m64n96k16");
675                }
676                Shape::M64n8k16 => {
677                    push_directive(tokens, "m64n8k16");
678                }
679            }
680            match &self.dtype {
681                Dtype::F32 => {
682                    push_directive(tokens, "f32");
683                }
684            }
685            push_directive(tokens, "bf16");
686            push_directive(tokens, "bf16");
687            if spaced {
688                tokens.push(PtxToken::Space);
689            }
690            self.d.unparse_tokens_mode(tokens, spaced);
691            tokens.push(PtxToken::Comma);
692            if spaced {
693                tokens.push(PtxToken::Space);
694            }
695            self.a.unparse_tokens_mode(tokens, spaced);
696            tokens.push(PtxToken::Comma);
697            if spaced {
698                tokens.push(PtxToken::Space);
699            }
700            self.b_desc.unparse_tokens_mode(tokens, spaced);
701            tokens.push(PtxToken::Comma);
702            if spaced {
703                tokens.push(PtxToken::Space);
704            }
705            self.scale_d.unparse_tokens_mode(tokens, spaced);
706            tokens.push(PtxToken::Comma);
707            if spaced {
708                tokens.push(PtxToken::Space);
709            }
710            self.imm_scale_a.unparse_tokens_mode(tokens, spaced);
711            tokens.push(PtxToken::Comma);
712            if spaced {
713                tokens.push(PtxToken::Space);
714            }
715            self.imm_scale_b.unparse_tokens_mode(tokens, spaced);
716            tokens.push(PtxToken::Comma);
717            if spaced {
718                tokens.push(PtxToken::Space);
719            }
720            self.imm_trans_b.unparse_tokens_mode(tokens, spaced);
721            tokens.push(PtxToken::Semicolon);
722            if spaced {
723                tokens.push(PtxToken::Newline);
724            }
725        }
726    }
727}
728
729pub mod section_2 {
730    use super::*;
731    use crate::r#type::instruction::wgmma_mma_async::section_2::*;
732
733    impl PtxUnparser for WgmmaMmaAsyncSyncAlignedShapeDtypeTf32Tf32 {
734        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
735            self.unparse_tokens_mode(tokens, false);
736        }
737        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
738            push_opcode(tokens, "wgmma");
739            push_directive(tokens, "mma_async");
740            push_directive(tokens, "sync");
741            push_directive(tokens, "aligned");
742            match &self.shape {
743                Shape::M64n104k8 => {
744                    push_directive(tokens, "m64n104k8");
745                }
746                Shape::M64n112k8 => {
747                    push_directive(tokens, "m64n112k8");
748                }
749                Shape::M64n120k8 => {
750                    push_directive(tokens, "m64n120k8");
751                }
752                Shape::M64n128k8 => {
753                    push_directive(tokens, "m64n128k8");
754                }
755                Shape::M64n136k8 => {
756                    push_directive(tokens, "m64n136k8");
757                }
758                Shape::M64n144k8 => {
759                    push_directive(tokens, "m64n144k8");
760                }
761                Shape::M64n152k8 => {
762                    push_directive(tokens, "m64n152k8");
763                }
764                Shape::M64n160k8 => {
765                    push_directive(tokens, "m64n160k8");
766                }
767                Shape::M64n168k8 => {
768                    push_directive(tokens, "m64n168k8");
769                }
770                Shape::M64n176k8 => {
771                    push_directive(tokens, "m64n176k8");
772                }
773                Shape::M64n184k8 => {
774                    push_directive(tokens, "m64n184k8");
775                }
776                Shape::M64n192k8 => {
777                    push_directive(tokens, "m64n192k8");
778                }
779                Shape::M64n200k8 => {
780                    push_directive(tokens, "m64n200k8");
781                }
782                Shape::M64n208k8 => {
783                    push_directive(tokens, "m64n208k8");
784                }
785                Shape::M64n216k8 => {
786                    push_directive(tokens, "m64n216k8");
787                }
788                Shape::M64n224k8 => {
789                    push_directive(tokens, "m64n224k8");
790                }
791                Shape::M64n232k8 => {
792                    push_directive(tokens, "m64n232k8");
793                }
794                Shape::M64n240k8 => {
795                    push_directive(tokens, "m64n240k8");
796                }
797                Shape::M64n248k8 => {
798                    push_directive(tokens, "m64n248k8");
799                }
800                Shape::M64n256k8 => {
801                    push_directive(tokens, "m64n256k8");
802                }
803                Shape::M64n16k8 => {
804                    push_directive(tokens, "m64n16k8");
805                }
806                Shape::M64n24k8 => {
807                    push_directive(tokens, "m64n24k8");
808                }
809                Shape::M64n32k8 => {
810                    push_directive(tokens, "m64n32k8");
811                }
812                Shape::M64n40k8 => {
813                    push_directive(tokens, "m64n40k8");
814                }
815                Shape::M64n48k8 => {
816                    push_directive(tokens, "m64n48k8");
817                }
818                Shape::M64n56k8 => {
819                    push_directive(tokens, "m64n56k8");
820                }
821                Shape::M64n64k8 => {
822                    push_directive(tokens, "m64n64k8");
823                }
824                Shape::M64n72k8 => {
825                    push_directive(tokens, "m64n72k8");
826                }
827                Shape::M64n80k8 => {
828                    push_directive(tokens, "m64n80k8");
829                }
830                Shape::M64n88k8 => {
831                    push_directive(tokens, "m64n88k8");
832                }
833                Shape::M64n96k8 => {
834                    push_directive(tokens, "m64n96k8");
835                }
836                Shape::M64n8k8 => {
837                    push_directive(tokens, "m64n8k8");
838                }
839            }
840            match &self.dtype {
841                Dtype::F32 => {
842                    push_directive(tokens, "f32");
843                }
844            }
845            push_directive(tokens, "tf32");
846            push_directive(tokens, "tf32");
847            if spaced {
848                tokens.push(PtxToken::Space);
849            }
850            self.d.unparse_tokens_mode(tokens, spaced);
851            tokens.push(PtxToken::Comma);
852            if spaced {
853                tokens.push(PtxToken::Space);
854            }
855            self.a_desc.unparse_tokens_mode(tokens, spaced);
856            tokens.push(PtxToken::Comma);
857            if spaced {
858                tokens.push(PtxToken::Space);
859            }
860            self.b_desc.unparse_tokens_mode(tokens, spaced);
861            tokens.push(PtxToken::Comma);
862            if spaced {
863                tokens.push(PtxToken::Space);
864            }
865            self.scale_d.unparse_tokens_mode(tokens, spaced);
866            tokens.push(PtxToken::Comma);
867            if spaced {
868                tokens.push(PtxToken::Space);
869            }
870            self.imm_scale_a.unparse_tokens_mode(tokens, spaced);
871            tokens.push(PtxToken::Comma);
872            if spaced {
873                tokens.push(PtxToken::Space);
874            }
875            self.imm_scale_b.unparse_tokens_mode(tokens, spaced);
876            tokens.push(PtxToken::Semicolon);
877            if spaced {
878                tokens.push(PtxToken::Newline);
879            }
880        }
881    }
882
883    impl PtxUnparser for WgmmaMmaAsyncSyncAlignedShapeDtypeTf32Tf321 {
884        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
885            self.unparse_tokens_mode(tokens, false);
886        }
887        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
888            push_opcode(tokens, "wgmma");
889            push_directive(tokens, "mma_async");
890            push_directive(tokens, "sync");
891            push_directive(tokens, "aligned");
892            match &self.shape {
893                Shape::M64n104k8 => {
894                    push_directive(tokens, "m64n104k8");
895                }
896                Shape::M64n112k8 => {
897                    push_directive(tokens, "m64n112k8");
898                }
899                Shape::M64n120k8 => {
900                    push_directive(tokens, "m64n120k8");
901                }
902                Shape::M64n128k8 => {
903                    push_directive(tokens, "m64n128k8");
904                }
905                Shape::M64n136k8 => {
906                    push_directive(tokens, "m64n136k8");
907                }
908                Shape::M64n144k8 => {
909                    push_directive(tokens, "m64n144k8");
910                }
911                Shape::M64n152k8 => {
912                    push_directive(tokens, "m64n152k8");
913                }
914                Shape::M64n160k8 => {
915                    push_directive(tokens, "m64n160k8");
916                }
917                Shape::M64n168k8 => {
918                    push_directive(tokens, "m64n168k8");
919                }
920                Shape::M64n176k8 => {
921                    push_directive(tokens, "m64n176k8");
922                }
923                Shape::M64n184k8 => {
924                    push_directive(tokens, "m64n184k8");
925                }
926                Shape::M64n192k8 => {
927                    push_directive(tokens, "m64n192k8");
928                }
929                Shape::M64n200k8 => {
930                    push_directive(tokens, "m64n200k8");
931                }
932                Shape::M64n208k8 => {
933                    push_directive(tokens, "m64n208k8");
934                }
935                Shape::M64n216k8 => {
936                    push_directive(tokens, "m64n216k8");
937                }
938                Shape::M64n224k8 => {
939                    push_directive(tokens, "m64n224k8");
940                }
941                Shape::M64n232k8 => {
942                    push_directive(tokens, "m64n232k8");
943                }
944                Shape::M64n240k8 => {
945                    push_directive(tokens, "m64n240k8");
946                }
947                Shape::M64n248k8 => {
948                    push_directive(tokens, "m64n248k8");
949                }
950                Shape::M64n256k8 => {
951                    push_directive(tokens, "m64n256k8");
952                }
953                Shape::M64n16k8 => {
954                    push_directive(tokens, "m64n16k8");
955                }
956                Shape::M64n24k8 => {
957                    push_directive(tokens, "m64n24k8");
958                }
959                Shape::M64n32k8 => {
960                    push_directive(tokens, "m64n32k8");
961                }
962                Shape::M64n40k8 => {
963                    push_directive(tokens, "m64n40k8");
964                }
965                Shape::M64n48k8 => {
966                    push_directive(tokens, "m64n48k8");
967                }
968                Shape::M64n56k8 => {
969                    push_directive(tokens, "m64n56k8");
970                }
971                Shape::M64n64k8 => {
972                    push_directive(tokens, "m64n64k8");
973                }
974                Shape::M64n72k8 => {
975                    push_directive(tokens, "m64n72k8");
976                }
977                Shape::M64n80k8 => {
978                    push_directive(tokens, "m64n80k8");
979                }
980                Shape::M64n88k8 => {
981                    push_directive(tokens, "m64n88k8");
982                }
983                Shape::M64n96k8 => {
984                    push_directive(tokens, "m64n96k8");
985                }
986                Shape::M64n8k8 => {
987                    push_directive(tokens, "m64n8k8");
988                }
989            }
990            match &self.dtype {
991                Dtype::F32 => {
992                    push_directive(tokens, "f32");
993                }
994            }
995            push_directive(tokens, "tf32");
996            push_directive(tokens, "tf32");
997            if spaced {
998                tokens.push(PtxToken::Space);
999            }
1000            self.d.unparse_tokens_mode(tokens, spaced);
1001            tokens.push(PtxToken::Comma);
1002            if spaced {
1003                tokens.push(PtxToken::Space);
1004            }
1005            self.a.unparse_tokens_mode(tokens, spaced);
1006            tokens.push(PtxToken::Comma);
1007            if spaced {
1008                tokens.push(PtxToken::Space);
1009            }
1010            self.b_desc.unparse_tokens_mode(tokens, spaced);
1011            tokens.push(PtxToken::Comma);
1012            if spaced {
1013                tokens.push(PtxToken::Space);
1014            }
1015            self.scale_d.unparse_tokens_mode(tokens, spaced);
1016            tokens.push(PtxToken::Comma);
1017            if spaced {
1018                tokens.push(PtxToken::Space);
1019            }
1020            self.imm_scale_a.unparse_tokens_mode(tokens, spaced);
1021            tokens.push(PtxToken::Comma);
1022            if spaced {
1023                tokens.push(PtxToken::Space);
1024            }
1025            self.imm_scale_b.unparse_tokens_mode(tokens, spaced);
1026            tokens.push(PtxToken::Semicolon);
1027            if spaced {
1028                tokens.push(PtxToken::Newline);
1029            }
1030        }
1031    }
1032}
1033
1034pub mod section_3 {
1035    use super::*;
1036    use crate::r#type::instruction::wgmma_mma_async::section_3::*;
1037
1038    impl PtxUnparser for WgmmaMmaAsyncSyncAlignedShapeDtypeAtypeBtype {
1039        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
1040            self.unparse_tokens_mode(tokens, false);
1041        }
1042        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
1043            push_opcode(tokens, "wgmma");
1044            push_directive(tokens, "mma_async");
1045            push_directive(tokens, "sync");
1046            push_directive(tokens, "aligned");
1047            match &self.shape {
1048                Shape::M64n104k32 => {
1049                    push_directive(tokens, "m64n104k32");
1050                }
1051                Shape::M64n112k32 => {
1052                    push_directive(tokens, "m64n112k32");
1053                }
1054                Shape::M64n120k32 => {
1055                    push_directive(tokens, "m64n120k32");
1056                }
1057                Shape::M64n128k32 => {
1058                    push_directive(tokens, "m64n128k32");
1059                }
1060                Shape::M64n136k32 => {
1061                    push_directive(tokens, "m64n136k32");
1062                }
1063                Shape::M64n144k32 => {
1064                    push_directive(tokens, "m64n144k32");
1065                }
1066                Shape::M64n152k32 => {
1067                    push_directive(tokens, "m64n152k32");
1068                }
1069                Shape::M64n160k32 => {
1070                    push_directive(tokens, "m64n160k32");
1071                }
1072                Shape::M64n168k32 => {
1073                    push_directive(tokens, "m64n168k32");
1074                }
1075                Shape::M64n176k32 => {
1076                    push_directive(tokens, "m64n176k32");
1077                }
1078                Shape::M64n184k32 => {
1079                    push_directive(tokens, "m64n184k32");
1080                }
1081                Shape::M64n192k32 => {
1082                    push_directive(tokens, "m64n192k32");
1083                }
1084                Shape::M64n200k32 => {
1085                    push_directive(tokens, "m64n200k32");
1086                }
1087                Shape::M64n208k32 => {
1088                    push_directive(tokens, "m64n208k32");
1089                }
1090                Shape::M64n216k32 => {
1091                    push_directive(tokens, "m64n216k32");
1092                }
1093                Shape::M64n224k32 => {
1094                    push_directive(tokens, "m64n224k32");
1095                }
1096                Shape::M64n232k32 => {
1097                    push_directive(tokens, "m64n232k32");
1098                }
1099                Shape::M64n240k32 => {
1100                    push_directive(tokens, "m64n240k32");
1101                }
1102                Shape::M64n248k32 => {
1103                    push_directive(tokens, "m64n248k32");
1104                }
1105                Shape::M64n256k32 => {
1106                    push_directive(tokens, "m64n256k32");
1107                }
1108                Shape::M64n16k32 => {
1109                    push_directive(tokens, "m64n16k32");
1110                }
1111                Shape::M64n24k32 => {
1112                    push_directive(tokens, "m64n24k32");
1113                }
1114                Shape::M64n32k32 => {
1115                    push_directive(tokens, "m64n32k32");
1116                }
1117                Shape::M64n40k32 => {
1118                    push_directive(tokens, "m64n40k32");
1119                }
1120                Shape::M64n48k32 => {
1121                    push_directive(tokens, "m64n48k32");
1122                }
1123                Shape::M64n56k32 => {
1124                    push_directive(tokens, "m64n56k32");
1125                }
1126                Shape::M64n64k32 => {
1127                    push_directive(tokens, "m64n64k32");
1128                }
1129                Shape::M64n72k32 => {
1130                    push_directive(tokens, "m64n72k32");
1131                }
1132                Shape::M64n80k32 => {
1133                    push_directive(tokens, "m64n80k32");
1134                }
1135                Shape::M64n88k32 => {
1136                    push_directive(tokens, "m64n88k32");
1137                }
1138                Shape::M64n96k32 => {
1139                    push_directive(tokens, "m64n96k32");
1140                }
1141                Shape::M64n8k32 => {
1142                    push_directive(tokens, "m64n8k32");
1143                }
1144            }
1145            match &self.dtype {
1146                Dtype::F16 => {
1147                    push_directive(tokens, "f16");
1148                }
1149                Dtype::F32 => {
1150                    push_directive(tokens, "f32");
1151                }
1152            }
1153            match &self.atype {
1154                Atype::E4m3 => {
1155                    push_directive(tokens, "e4m3");
1156                }
1157                Atype::E5m2 => {
1158                    push_directive(tokens, "e5m2");
1159                }
1160            }
1161            match &self.btype {
1162                Btype::E4m3 => {
1163                    push_directive(tokens, "e4m3");
1164                }
1165                Btype::E5m2 => {
1166                    push_directive(tokens, "e5m2");
1167                }
1168            }
1169            if spaced {
1170                tokens.push(PtxToken::Space);
1171            }
1172            self.d.unparse_tokens_mode(tokens, spaced);
1173            tokens.push(PtxToken::Comma);
1174            if spaced {
1175                tokens.push(PtxToken::Space);
1176            }
1177            self.a_desc.unparse_tokens_mode(tokens, spaced);
1178            tokens.push(PtxToken::Comma);
1179            if spaced {
1180                tokens.push(PtxToken::Space);
1181            }
1182            self.b_desc.unparse_tokens_mode(tokens, spaced);
1183            tokens.push(PtxToken::Comma);
1184            if spaced {
1185                tokens.push(PtxToken::Space);
1186            }
1187            self.scale_d.unparse_tokens_mode(tokens, spaced);
1188            tokens.push(PtxToken::Comma);
1189            if spaced {
1190                tokens.push(PtxToken::Space);
1191            }
1192            self.imm_scale_a.unparse_tokens_mode(tokens, spaced);
1193            tokens.push(PtxToken::Comma);
1194            if spaced {
1195                tokens.push(PtxToken::Space);
1196            }
1197            self.imm_scale_b.unparse_tokens_mode(tokens, spaced);
1198            tokens.push(PtxToken::Semicolon);
1199            if spaced {
1200                tokens.push(PtxToken::Newline);
1201            }
1202        }
1203    }
1204
1205    impl PtxUnparser for WgmmaMmaAsyncSyncAlignedShapeDtypeAtypeBtype1 {
1206        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
1207            self.unparse_tokens_mode(tokens, false);
1208        }
1209        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
1210            push_opcode(tokens, "wgmma");
1211            push_directive(tokens, "mma_async");
1212            push_directive(tokens, "sync");
1213            push_directive(tokens, "aligned");
1214            match &self.shape {
1215                Shape::M64n104k32 => {
1216                    push_directive(tokens, "m64n104k32");
1217                }
1218                Shape::M64n112k32 => {
1219                    push_directive(tokens, "m64n112k32");
1220                }
1221                Shape::M64n120k32 => {
1222                    push_directive(tokens, "m64n120k32");
1223                }
1224                Shape::M64n128k32 => {
1225                    push_directive(tokens, "m64n128k32");
1226                }
1227                Shape::M64n136k32 => {
1228                    push_directive(tokens, "m64n136k32");
1229                }
1230                Shape::M64n144k32 => {
1231                    push_directive(tokens, "m64n144k32");
1232                }
1233                Shape::M64n152k32 => {
1234                    push_directive(tokens, "m64n152k32");
1235                }
1236                Shape::M64n160k32 => {
1237                    push_directive(tokens, "m64n160k32");
1238                }
1239                Shape::M64n168k32 => {
1240                    push_directive(tokens, "m64n168k32");
1241                }
1242                Shape::M64n176k32 => {
1243                    push_directive(tokens, "m64n176k32");
1244                }
1245                Shape::M64n184k32 => {
1246                    push_directive(tokens, "m64n184k32");
1247                }
1248                Shape::M64n192k32 => {
1249                    push_directive(tokens, "m64n192k32");
1250                }
1251                Shape::M64n200k32 => {
1252                    push_directive(tokens, "m64n200k32");
1253                }
1254                Shape::M64n208k32 => {
1255                    push_directive(tokens, "m64n208k32");
1256                }
1257                Shape::M64n216k32 => {
1258                    push_directive(tokens, "m64n216k32");
1259                }
1260                Shape::M64n224k32 => {
1261                    push_directive(tokens, "m64n224k32");
1262                }
1263                Shape::M64n232k32 => {
1264                    push_directive(tokens, "m64n232k32");
1265                }
1266                Shape::M64n240k32 => {
1267                    push_directive(tokens, "m64n240k32");
1268                }
1269                Shape::M64n248k32 => {
1270                    push_directive(tokens, "m64n248k32");
1271                }
1272                Shape::M64n256k32 => {
1273                    push_directive(tokens, "m64n256k32");
1274                }
1275                Shape::M64n16k32 => {
1276                    push_directive(tokens, "m64n16k32");
1277                }
1278                Shape::M64n24k32 => {
1279                    push_directive(tokens, "m64n24k32");
1280                }
1281                Shape::M64n32k32 => {
1282                    push_directive(tokens, "m64n32k32");
1283                }
1284                Shape::M64n40k32 => {
1285                    push_directive(tokens, "m64n40k32");
1286                }
1287                Shape::M64n48k32 => {
1288                    push_directive(tokens, "m64n48k32");
1289                }
1290                Shape::M64n56k32 => {
1291                    push_directive(tokens, "m64n56k32");
1292                }
1293                Shape::M64n64k32 => {
1294                    push_directive(tokens, "m64n64k32");
1295                }
1296                Shape::M64n72k32 => {
1297                    push_directive(tokens, "m64n72k32");
1298                }
1299                Shape::M64n80k32 => {
1300                    push_directive(tokens, "m64n80k32");
1301                }
1302                Shape::M64n88k32 => {
1303                    push_directive(tokens, "m64n88k32");
1304                }
1305                Shape::M64n96k32 => {
1306                    push_directive(tokens, "m64n96k32");
1307                }
1308                Shape::M64n8k32 => {
1309                    push_directive(tokens, "m64n8k32");
1310                }
1311            }
1312            match &self.dtype {
1313                Dtype::F16 => {
1314                    push_directive(tokens, "f16");
1315                }
1316                Dtype::F32 => {
1317                    push_directive(tokens, "f32");
1318                }
1319            }
1320            match &self.atype {
1321                Atype::E4m3 => {
1322                    push_directive(tokens, "e4m3");
1323                }
1324                Atype::E5m2 => {
1325                    push_directive(tokens, "e5m2");
1326                }
1327            }
1328            match &self.btype {
1329                Btype::E4m3 => {
1330                    push_directive(tokens, "e4m3");
1331                }
1332                Btype::E5m2 => {
1333                    push_directive(tokens, "e5m2");
1334                }
1335            }
1336            if spaced {
1337                tokens.push(PtxToken::Space);
1338            }
1339            self.d.unparse_tokens_mode(tokens, spaced);
1340            tokens.push(PtxToken::Comma);
1341            if spaced {
1342                tokens.push(PtxToken::Space);
1343            }
1344            self.a.unparse_tokens_mode(tokens, spaced);
1345            tokens.push(PtxToken::Comma);
1346            if spaced {
1347                tokens.push(PtxToken::Space);
1348            }
1349            self.b_desc.unparse_tokens_mode(tokens, spaced);
1350            tokens.push(PtxToken::Comma);
1351            if spaced {
1352                tokens.push(PtxToken::Space);
1353            }
1354            self.scale_d.unparse_tokens_mode(tokens, spaced);
1355            tokens.push(PtxToken::Comma);
1356            if spaced {
1357                tokens.push(PtxToken::Space);
1358            }
1359            self.imm_scale_a.unparse_tokens_mode(tokens, spaced);
1360            tokens.push(PtxToken::Comma);
1361            if spaced {
1362                tokens.push(PtxToken::Space);
1363            }
1364            self.imm_scale_b.unparse_tokens_mode(tokens, spaced);
1365            tokens.push(PtxToken::Semicolon);
1366            if spaced {
1367                tokens.push(PtxToken::Newline);
1368            }
1369        }
1370    }
1371}
1372
1373pub mod section_4 {
1374    use super::*;
1375    use crate::r#type::instruction::wgmma_mma_async::section_4::*;
1376
1377    impl PtxUnparser for WgmmaMmaAsyncSyncAlignedShapeSatfiniteS32AtypeBtype {
1378        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
1379            self.unparse_tokens_mode(tokens, false);
1380        }
1381        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
1382            push_opcode(tokens, "wgmma");
1383            push_directive(tokens, "mma_async");
1384            push_directive(tokens, "sync");
1385            push_directive(tokens, "aligned");
1386            match &self.shape {
1387                Shape::M64n112k32 => {
1388                    push_directive(tokens, "m64n112k32");
1389                }
1390                Shape::M64n128k32 => {
1391                    push_directive(tokens, "m64n128k32");
1392                }
1393                Shape::M64n144k32 => {
1394                    push_directive(tokens, "m64n144k32");
1395                }
1396                Shape::M64n160k32 => {
1397                    push_directive(tokens, "m64n160k32");
1398                }
1399                Shape::M64n176k32 => {
1400                    push_directive(tokens, "m64n176k32");
1401                }
1402                Shape::M64n192k32 => {
1403                    push_directive(tokens, "m64n192k32");
1404                }
1405                Shape::M64n208k32 => {
1406                    push_directive(tokens, "m64n208k32");
1407                }
1408                Shape::M64n224k32 => {
1409                    push_directive(tokens, "m64n224k32");
1410                }
1411                Shape::M64n16k32 => {
1412                    push_directive(tokens, "m64n16k32");
1413                }
1414                Shape::M64n24k32 => {
1415                    push_directive(tokens, "m64n24k32");
1416                }
1417                Shape::M64n32k32 => {
1418                    push_directive(tokens, "m64n32k32");
1419                }
1420                Shape::M64n48k32 => {
1421                    push_directive(tokens, "m64n48k32");
1422                }
1423                Shape::M64n64k32 => {
1424                    push_directive(tokens, "m64n64k32");
1425                }
1426                Shape::M64n80k32 => {
1427                    push_directive(tokens, "m64n80k32");
1428                }
1429                Shape::M64n96k32 => {
1430                    push_directive(tokens, "m64n96k32");
1431                }
1432                Shape::M64n8k32 => {
1433                    push_directive(tokens, "m64n8k32");
1434                }
1435            }
1436            if self.satfinite {
1437                push_directive(tokens, "satfinite");
1438            }
1439            push_directive(tokens, "s32");
1440            match &self.atype {
1441                Atype::S8 => {
1442                    push_directive(tokens, "s8");
1443                }
1444                Atype::U8 => {
1445                    push_directive(tokens, "u8");
1446                }
1447            }
1448            match &self.btype {
1449                Btype::S8 => {
1450                    push_directive(tokens, "s8");
1451                }
1452                Btype::U8 => {
1453                    push_directive(tokens, "u8");
1454                }
1455            }
1456            if spaced {
1457                tokens.push(PtxToken::Space);
1458            }
1459            self.d.unparse_tokens_mode(tokens, spaced);
1460            tokens.push(PtxToken::Comma);
1461            if spaced {
1462                tokens.push(PtxToken::Space);
1463            }
1464            self.a_desc.unparse_tokens_mode(tokens, spaced);
1465            tokens.push(PtxToken::Comma);
1466            if spaced {
1467                tokens.push(PtxToken::Space);
1468            }
1469            self.b_desc.unparse_tokens_mode(tokens, spaced);
1470            tokens.push(PtxToken::Comma);
1471            if spaced {
1472                tokens.push(PtxToken::Space);
1473            }
1474            self.scale_d.unparse_tokens_mode(tokens, spaced);
1475            tokens.push(PtxToken::Semicolon);
1476            if spaced {
1477                tokens.push(PtxToken::Newline);
1478            }
1479        }
1480    }
1481
1482    impl PtxUnparser for WgmmaMmaAsyncSyncAlignedShapeSatfiniteS32AtypeBtype1 {
1483        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
1484            self.unparse_tokens_mode(tokens, false);
1485        }
1486        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
1487            push_opcode(tokens, "wgmma");
1488            push_directive(tokens, "mma_async");
1489            push_directive(tokens, "sync");
1490            push_directive(tokens, "aligned");
1491            match &self.shape {
1492                Shape::M64n112k32 => {
1493                    push_directive(tokens, "m64n112k32");
1494                }
1495                Shape::M64n128k32 => {
1496                    push_directive(tokens, "m64n128k32");
1497                }
1498                Shape::M64n144k32 => {
1499                    push_directive(tokens, "m64n144k32");
1500                }
1501                Shape::M64n160k32 => {
1502                    push_directive(tokens, "m64n160k32");
1503                }
1504                Shape::M64n176k32 => {
1505                    push_directive(tokens, "m64n176k32");
1506                }
1507                Shape::M64n192k32 => {
1508                    push_directive(tokens, "m64n192k32");
1509                }
1510                Shape::M64n208k32 => {
1511                    push_directive(tokens, "m64n208k32");
1512                }
1513                Shape::M64n224k32 => {
1514                    push_directive(tokens, "m64n224k32");
1515                }
1516                Shape::M64n16k32 => {
1517                    push_directive(tokens, "m64n16k32");
1518                }
1519                Shape::M64n24k32 => {
1520                    push_directive(tokens, "m64n24k32");
1521                }
1522                Shape::M64n32k32 => {
1523                    push_directive(tokens, "m64n32k32");
1524                }
1525                Shape::M64n48k32 => {
1526                    push_directive(tokens, "m64n48k32");
1527                }
1528                Shape::M64n64k32 => {
1529                    push_directive(tokens, "m64n64k32");
1530                }
1531                Shape::M64n80k32 => {
1532                    push_directive(tokens, "m64n80k32");
1533                }
1534                Shape::M64n96k32 => {
1535                    push_directive(tokens, "m64n96k32");
1536                }
1537                Shape::M64n8k32 => {
1538                    push_directive(tokens, "m64n8k32");
1539                }
1540            }
1541            if self.satfinite {
1542                push_directive(tokens, "satfinite");
1543            }
1544            push_directive(tokens, "s32");
1545            match &self.atype {
1546                Atype::S8 => {
1547                    push_directive(tokens, "s8");
1548                }
1549                Atype::U8 => {
1550                    push_directive(tokens, "u8");
1551                }
1552            }
1553            match &self.btype {
1554                Btype::S8 => {
1555                    push_directive(tokens, "s8");
1556                }
1557                Btype::U8 => {
1558                    push_directive(tokens, "u8");
1559                }
1560            }
1561            if spaced {
1562                tokens.push(PtxToken::Space);
1563            }
1564            self.d.unparse_tokens_mode(tokens, spaced);
1565            tokens.push(PtxToken::Comma);
1566            if spaced {
1567                tokens.push(PtxToken::Space);
1568            }
1569            self.a.unparse_tokens_mode(tokens, spaced);
1570            tokens.push(PtxToken::Comma);
1571            if spaced {
1572                tokens.push(PtxToken::Space);
1573            }
1574            self.b_desc.unparse_tokens_mode(tokens, spaced);
1575            tokens.push(PtxToken::Comma);
1576            if spaced {
1577                tokens.push(PtxToken::Space);
1578            }
1579            self.scale_d.unparse_tokens_mode(tokens, spaced);
1580            tokens.push(PtxToken::Semicolon);
1581            if spaced {
1582                tokens.push(PtxToken::Newline);
1583            }
1584        }
1585    }
1586}
1587
1588pub mod section_5 {
1589    use super::*;
1590    use crate::r#type::instruction::wgmma_mma_async::section_5::*;
1591
1592    impl PtxUnparser for WgmmaMmaAsyncSyncAlignedShapeS32B1B1OpPopc {
1593        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
1594            self.unparse_tokens_mode(tokens, false);
1595        }
1596        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
1597            push_opcode(tokens, "wgmma");
1598            push_directive(tokens, "mma_async");
1599            push_directive(tokens, "sync");
1600            push_directive(tokens, "aligned");
1601            match &self.shape {
1602                Shape::M64n112k256 => {
1603                    push_directive(tokens, "m64n112k256");
1604                }
1605                Shape::M64n128k256 => {
1606                    push_directive(tokens, "m64n128k256");
1607                }
1608                Shape::M64n144k256 => {
1609                    push_directive(tokens, "m64n144k256");
1610                }
1611                Shape::M64n160k256 => {
1612                    push_directive(tokens, "m64n160k256");
1613                }
1614                Shape::M64n176k256 => {
1615                    push_directive(tokens, "m64n176k256");
1616                }
1617                Shape::M64n192k256 => {
1618                    push_directive(tokens, "m64n192k256");
1619                }
1620                Shape::M64n208k256 => {
1621                    push_directive(tokens, "m64n208k256");
1622                }
1623                Shape::M64n224k256 => {
1624                    push_directive(tokens, "m64n224k256");
1625                }
1626                Shape::M64n240k256 => {
1627                    push_directive(tokens, "m64n240k256");
1628                }
1629                Shape::M64n256k256 => {
1630                    push_directive(tokens, "m64n256k256");
1631                }
1632                Shape::M64n16k256 => {
1633                    push_directive(tokens, "m64n16k256");
1634                }
1635                Shape::M64n24k256 => {
1636                    push_directive(tokens, "m64n24k256");
1637                }
1638                Shape::M64n32k256 => {
1639                    push_directive(tokens, "m64n32k256");
1640                }
1641                Shape::M64n48k256 => {
1642                    push_directive(tokens, "m64n48k256");
1643                }
1644                Shape::M64n64k256 => {
1645                    push_directive(tokens, "m64n64k256");
1646                }
1647                Shape::M64n80k256 => {
1648                    push_directive(tokens, "m64n80k256");
1649                }
1650                Shape::M64n96k256 => {
1651                    push_directive(tokens, "m64n96k256");
1652                }
1653                Shape::M64n8k256 => {
1654                    push_directive(tokens, "m64n8k256");
1655                }
1656            }
1657            push_directive(tokens, "s32");
1658            push_directive(tokens, "b1");
1659            push_directive(tokens, "b1");
1660            match &self.op {
1661                Op::And => {
1662                    push_directive(tokens, "and");
1663                }
1664            }
1665            push_directive(tokens, "popc");
1666            if spaced {
1667                tokens.push(PtxToken::Space);
1668            }
1669            self.d.unparse_tokens_mode(tokens, spaced);
1670            tokens.push(PtxToken::Comma);
1671            if spaced {
1672                tokens.push(PtxToken::Space);
1673            }
1674            self.a_desc.unparse_tokens_mode(tokens, spaced);
1675            tokens.push(PtxToken::Comma);
1676            if spaced {
1677                tokens.push(PtxToken::Space);
1678            }
1679            self.b_desc.unparse_tokens_mode(tokens, spaced);
1680            tokens.push(PtxToken::Comma);
1681            if spaced {
1682                tokens.push(PtxToken::Space);
1683            }
1684            self.scale_d.unparse_tokens_mode(tokens, spaced);
1685            tokens.push(PtxToken::Semicolon);
1686            if spaced {
1687                tokens.push(PtxToken::Newline);
1688            }
1689        }
1690    }
1691
1692    impl PtxUnparser for WgmmaMmaAsyncSyncAlignedShapeS32B1B1OpPopc1 {
1693        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
1694            self.unparse_tokens_mode(tokens, false);
1695        }
1696        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
1697            push_opcode(tokens, "wgmma");
1698            push_directive(tokens, "mma_async");
1699            push_directive(tokens, "sync");
1700            push_directive(tokens, "aligned");
1701            match &self.shape {
1702                Shape::M64n112k256 => {
1703                    push_directive(tokens, "m64n112k256");
1704                }
1705                Shape::M64n128k256 => {
1706                    push_directive(tokens, "m64n128k256");
1707                }
1708                Shape::M64n144k256 => {
1709                    push_directive(tokens, "m64n144k256");
1710                }
1711                Shape::M64n160k256 => {
1712                    push_directive(tokens, "m64n160k256");
1713                }
1714                Shape::M64n176k256 => {
1715                    push_directive(tokens, "m64n176k256");
1716                }
1717                Shape::M64n192k256 => {
1718                    push_directive(tokens, "m64n192k256");
1719                }
1720                Shape::M64n208k256 => {
1721                    push_directive(tokens, "m64n208k256");
1722                }
1723                Shape::M64n224k256 => {
1724                    push_directive(tokens, "m64n224k256");
1725                }
1726                Shape::M64n240k256 => {
1727                    push_directive(tokens, "m64n240k256");
1728                }
1729                Shape::M64n256k256 => {
1730                    push_directive(tokens, "m64n256k256");
1731                }
1732                Shape::M64n16k256 => {
1733                    push_directive(tokens, "m64n16k256");
1734                }
1735                Shape::M64n24k256 => {
1736                    push_directive(tokens, "m64n24k256");
1737                }
1738                Shape::M64n32k256 => {
1739                    push_directive(tokens, "m64n32k256");
1740                }
1741                Shape::M64n48k256 => {
1742                    push_directive(tokens, "m64n48k256");
1743                }
1744                Shape::M64n64k256 => {
1745                    push_directive(tokens, "m64n64k256");
1746                }
1747                Shape::M64n80k256 => {
1748                    push_directive(tokens, "m64n80k256");
1749                }
1750                Shape::M64n96k256 => {
1751                    push_directive(tokens, "m64n96k256");
1752                }
1753                Shape::M64n8k256 => {
1754                    push_directive(tokens, "m64n8k256");
1755                }
1756            }
1757            push_directive(tokens, "s32");
1758            push_directive(tokens, "b1");
1759            push_directive(tokens, "b1");
1760            match &self.op {
1761                Op::And => {
1762                    push_directive(tokens, "and");
1763                }
1764            }
1765            push_directive(tokens, "popc");
1766            if spaced {
1767                tokens.push(PtxToken::Space);
1768            }
1769            self.d.unparse_tokens_mode(tokens, spaced);
1770            tokens.push(PtxToken::Comma);
1771            if spaced {
1772                tokens.push(PtxToken::Space);
1773            }
1774            self.a.unparse_tokens_mode(tokens, spaced);
1775            tokens.push(PtxToken::Comma);
1776            if spaced {
1777                tokens.push(PtxToken::Space);
1778            }
1779            self.b_desc.unparse_tokens_mode(tokens, spaced);
1780            tokens.push(PtxToken::Comma);
1781            if spaced {
1782                tokens.push(PtxToken::Space);
1783            }
1784            self.scale_d.unparse_tokens_mode(tokens, spaced);
1785            tokens.push(PtxToken::Semicolon);
1786            if spaced {
1787                tokens.push(PtxToken::Newline);
1788            }
1789        }
1790    }
1791}
ptx_parser/unparser/instruction/wgmma_mma_async.rs

ptx_parser/unparser/instruction/
wgmma_mma_async.rs
