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

ptx_parser/unparser/instruction/
wgmma_mma_async.rs
