ptx_parser/unparser/instruction/
wgmma_mma_async_sp.rs

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

ptx_parser/unparser/instruction/
wgmma_mma_async_sp.rs
