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