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

ptx_parser/unparser/instruction/
wgmma_mma_async_sp.rs
