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