ptx_parser/unparser/instruction/
tcgen05_mma.rs

1//! Original PTX specification:
2//!
3//! // 1. Floating-point type without block scaling:
4//! tcgen05.mma.cta_group.kind   [d-tmem],  a-desc,  b-desc, idesc
5//! {, disable-output-lane }, enable-input-d {, scale-input-d};
6//! tcgen05.mma.cta_group.kind   [d-tmem], [a-tmem], b-desc, idesc
7//! {, disable-output-lane }, enable-input-d {, scale-input-d};
8//! .kind      = { .kind::f16, .kind::tf32, .kind::f8f6f4 };
9//! .cta_group = { .cta_group::1, .cta_group::2 };
10//! ------------------------------------------------------------------
11//! // 2. Floating-point type with block scaling:
12//! tcgen05.mma.cta_group.kind.block_scale{.scale_vectorsize}
13//! [d-tmem],  a-desc,  b-desc, idesc,
14//! [scale-A-tmem], [scale-B-tmem], enable-input-d;
15//! tcgen05.mma.cta_group.kind.block_scale{.scale_vectorsize}
16//! [d-tmem], [a-tmem], b-desc, idesc,
17//! [scale-A-tmem], [scale-B-tmem], enable-input-d;
18//! .kind = { .kind::mxf8f6f4, .kind::mxf4, .kind::mxf4nvf4 };
19//! .cta_group      = { .cta_group::1,   .cta_group::2 };
20//! .scale_vectorsize = { .scale_vec::1X, .scale_vec::2X, .scale_vec::4X, .block16, .block32 };
21//! ------------------------------------------------------------------
22//! // 3. Convolution MMA for floating-point type without block scaling:
23//! tcgen05.mma.cta_group.kind.collector_usage [d-tmem],  a-desc,  b-desc, idesc
24//! {, disable-output-lane }, enable-input-d {, scale-input-d};
25//! tcgen05.mma.cta_group.kind{.ashift}.collector_usage [d-tmem], [a-tmem], b-desc, idesc
26//! {, disable-output-lane }, enable-input-d {, scale-input-d};
27//! tcgen05.mma.cta_group.kind.ashift{.collector_usage} [d-tmem], [a-tmem], b-desc, idesc
28//! {, disable-output-lane }, enable-input-d {, scale-input-d};
29//! .kind      = { .kind::f16, .kind::tf32, .kind::f8f6f4 };
30//! .cta_group = { .cta_group::1,   .cta_group::2 };
31//! .collector_usage = { .collector::buffer::op };
32//! ::buffer         = { ::a };
33//! ::op             = { ::fill, ::use, ::lastuse, ::discard* };
34//! ------------------------------------------------------------------
35//! // 4. Activation Stationary MMA for floating-point type with block scaling:
36//! tcgen05.mma.cta_group.kind.block_scale{.scale_vectorsize}.collector_usage
37//! [d-tmem],  a-desc,  b-desc, idesc,
38//! [scale-A-tmem], [scale-B-tmem], enable-input-d;
39//! tcgen05.mma.cta_group.kind.block_scale{.scale_vectorsize}.collector_usage
40//! [d-tmem], [a-tmem], b-desc, idesc,
41//! [scale-A-tmem], [scale-B-tmem], enable-input-d;
42//! .cta_group       = { .cta_group::1,   .cta_group::2 };
43//! .scale_vectorsize  = { .scale_vec::1X, .scale_vec::2X, .scale_vec::4X, .block16, .block32 };
44//! .kind            = { .kind::mxf8f6f4, .kind::mxf4, .kind::mxf4nvf4 };
45//! .collector_usage = { .collector::buffer::op };
46//! ::buffer         = { ::a };
47//! ::op             = { ::fill, ::use, ::lastuse, ::discard* };
48//! ------------------------------------------------------------------
49//! // 5. Integer type:
50//! tcgen05.mma.cta_group.kind::i8  [d-tmem],  a-desc,  b-desc, idesc
51//! {, disable-output-lane }, enable-input-d;
52//! tcgen05.mma.cta_group.kind::i8  [d-tmem], [a-tmem], b-desc, idesc
53//! {, disable-output-lane }, enable-input-d;
54//! .cta_group = { .cta_group::1,   .cta_group::2  };
55//! ------------------------------------------------------------------
56//! // 6. Convolution MMA for integer type:
57//! tcgen05.mma.cta_group.kind::i8.collector_usage          [d-tmem],  a-desc,  b-desc, idesc
58//! {, disable-output-lane }, enable-input-d;
59//! tcgen05.mma.cta_group.kind::i8.ashift{.collector_usage} [d-tmem], [a-tmem], b-desc, idesc
60//! {, disable-output-lane }, enable-input-d;
61//! tcgen05.mma.cta_group.kind::i8{.ashift}.collector_usage [d-tmem], [a-tmem], b-desc, idesc
62//! {, disable-output-lane }, enable-input-d;
63//! .cta_group       = { .cta_group::1,   .cta_group::2  };
64//! .collector_usage = { .collector::buffer::op };
65//! ::buffer         = { ::a };
66//! ::op             = { ::fill, ::use, ::lastuse, ::discard* };
67
68#![allow(unused)]
69
70use crate::lexer::PtxToken;
71use crate::unparser::{PtxUnparser, common::*};
72
73pub mod section_0 {
74    use super::*;
75    use crate::r#type::instruction::tcgen05_mma::section_0::*;
76
77    impl PtxUnparser for Tcgen05MmaCtaGroupKind {
78        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
79            self.unparse_tokens_mode(tokens, false);
80        }
81        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
82            push_opcode(tokens, "tcgen05");
83            push_directive(tokens, "mma");
84            match &self.cta_group {
85                CtaGroup::CtaGroup1 => {
86                    push_directive(tokens, "cta_group::1");
87                }
88                CtaGroup::CtaGroup2 => {
89                    push_directive(tokens, "cta_group::2");
90                }
91            }
92            match &self.kind {
93                Kind::KindF8f6f4 => {
94                    push_directive(tokens, "kind::f8f6f4");
95                }
96                Kind::KindTf32 => {
97                    push_directive(tokens, "kind::tf32");
98                }
99                Kind::KindF16 => {
100                    push_directive(tokens, "kind::f16");
101                }
102            }
103            if spaced {
104                tokens.push(PtxToken::Space);
105            }
106            self.d_tmem.unparse_tokens_mode(tokens, spaced);
107            tokens.push(PtxToken::Comma);
108            if spaced {
109                tokens.push(PtxToken::Space);
110            }
111            self.a_desc.unparse_tokens_mode(tokens, spaced);
112            tokens.push(PtxToken::Comma);
113            if spaced {
114                tokens.push(PtxToken::Space);
115            }
116            self.b_desc.unparse_tokens_mode(tokens, spaced);
117            tokens.push(PtxToken::Comma);
118            if spaced {
119                tokens.push(PtxToken::Space);
120            }
121            self.idesc.unparse_tokens_mode(tokens, spaced);
122            if self.disable_output_lane.is_some() {
123                tokens.push(PtxToken::Comma);
124            }
125            if let Some(opt_0) = self.disable_output_lane.as_ref() {
126                if spaced {
127                    tokens.push(PtxToken::Space);
128                }
129                opt_0.unparse_tokens_mode(tokens, spaced);
130            }
131            tokens.push(PtxToken::Comma);
132            if spaced {
133                tokens.push(PtxToken::Space);
134            }
135            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
136            if self.scale_input_d.is_some() {
137                tokens.push(PtxToken::Comma);
138            }
139            if let Some(opt_1) = self.scale_input_d.as_ref() {
140                if spaced {
141                    tokens.push(PtxToken::Space);
142                }
143                opt_1.unparse_tokens_mode(tokens, spaced);
144            }
145            tokens.push(PtxToken::Semicolon);
146            if spaced {
147                tokens.push(PtxToken::Newline);
148            }
149        }
150    }
151
152    impl PtxUnparser for Tcgen05MmaCtaGroupKind1 {
153        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
154            self.unparse_tokens_mode(tokens, false);
155        }
156        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
157            push_opcode(tokens, "tcgen05");
158            push_directive(tokens, "mma");
159            match &self.cta_group {
160                CtaGroup::CtaGroup1 => {
161                    push_directive(tokens, "cta_group::1");
162                }
163                CtaGroup::CtaGroup2 => {
164                    push_directive(tokens, "cta_group::2");
165                }
166            }
167            match &self.kind {
168                Kind::KindF8f6f4 => {
169                    push_directive(tokens, "kind::f8f6f4");
170                }
171                Kind::KindTf32 => {
172                    push_directive(tokens, "kind::tf32");
173                }
174                Kind::KindF16 => {
175                    push_directive(tokens, "kind::f16");
176                }
177            }
178            if spaced {
179                tokens.push(PtxToken::Space);
180            }
181            self.d_tmem.unparse_tokens_mode(tokens, spaced);
182            tokens.push(PtxToken::Comma);
183            if spaced {
184                tokens.push(PtxToken::Space);
185            }
186            self.a_tmem.unparse_tokens_mode(tokens, spaced);
187            tokens.push(PtxToken::Comma);
188            if spaced {
189                tokens.push(PtxToken::Space);
190            }
191            self.b_desc.unparse_tokens_mode(tokens, spaced);
192            tokens.push(PtxToken::Comma);
193            if spaced {
194                tokens.push(PtxToken::Space);
195            }
196            self.idesc.unparse_tokens_mode(tokens, spaced);
197            if self.disable_output_lane.is_some() {
198                tokens.push(PtxToken::Comma);
199            }
200            if let Some(opt_2) = self.disable_output_lane.as_ref() {
201                if spaced {
202                    tokens.push(PtxToken::Space);
203                }
204                opt_2.unparse_tokens_mode(tokens, spaced);
205            }
206            tokens.push(PtxToken::Comma);
207            if spaced {
208                tokens.push(PtxToken::Space);
209            }
210            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
211            if self.scale_input_d.is_some() {
212                tokens.push(PtxToken::Comma);
213            }
214            if let Some(opt_3) = self.scale_input_d.as_ref() {
215                if spaced {
216                    tokens.push(PtxToken::Space);
217                }
218                opt_3.unparse_tokens_mode(tokens, spaced);
219            }
220            tokens.push(PtxToken::Semicolon);
221            if spaced {
222                tokens.push(PtxToken::Newline);
223            }
224        }
225    }
226}
227
228pub mod section_1 {
229    use super::*;
230    use crate::r#type::instruction::tcgen05_mma::section_1::*;
231
232    impl PtxUnparser for Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsize {
233        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
234            self.unparse_tokens_mode(tokens, false);
235        }
236        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
237            push_opcode(tokens, "tcgen05");
238            push_directive(tokens, "mma");
239            match &self.cta_group {
240                CtaGroup::CtaGroup1 => {
241                    push_directive(tokens, "cta_group::1");
242                }
243                CtaGroup::CtaGroup2 => {
244                    push_directive(tokens, "cta_group::2");
245                }
246            }
247            match &self.kind {
248                Kind::KindMxf8f6f4 => {
249                    push_directive(tokens, "kind::mxf8f6f4");
250                }
251                Kind::KindMxf4nvf4 => {
252                    push_directive(tokens, "kind::mxf4nvf4");
253                }
254                Kind::KindMxf4 => {
255                    push_directive(tokens, "kind::mxf4");
256                }
257            }
258            push_directive(tokens, "block_scale");
259            if let Some(scale_vectorsize_4) = self.scale_vectorsize.as_ref() {
260                match scale_vectorsize_4 {
261                    ScaleVectorsize::ScaleVec1x => {
262                        push_directive(tokens, "scale_vec::1X");
263                    }
264                    ScaleVectorsize::ScaleVec2x => {
265                        push_directive(tokens, "scale_vec::2X");
266                    }
267                    ScaleVectorsize::ScaleVec4x => {
268                        push_directive(tokens, "scale_vec::4X");
269                    }
270                    ScaleVectorsize::Block16 => {
271                        push_directive(tokens, "block16");
272                    }
273                    ScaleVectorsize::Block32 => {
274                        push_directive(tokens, "block32");
275                    }
276                }
277            }
278            if spaced {
279                tokens.push(PtxToken::Space);
280            }
281            self.d_tmem.unparse_tokens_mode(tokens, spaced);
282            tokens.push(PtxToken::Comma);
283            if spaced {
284                tokens.push(PtxToken::Space);
285            }
286            self.a_desc.unparse_tokens_mode(tokens, spaced);
287            tokens.push(PtxToken::Comma);
288            if spaced {
289                tokens.push(PtxToken::Space);
290            }
291            self.b_desc.unparse_tokens_mode(tokens, spaced);
292            tokens.push(PtxToken::Comma);
293            if spaced {
294                tokens.push(PtxToken::Space);
295            }
296            self.idesc.unparse_tokens_mode(tokens, spaced);
297            tokens.push(PtxToken::Comma);
298            if spaced {
299                tokens.push(PtxToken::Space);
300            }
301            self.scale_a_tmem.unparse_tokens_mode(tokens, spaced);
302            tokens.push(PtxToken::Comma);
303            if spaced {
304                tokens.push(PtxToken::Space);
305            }
306            self.scale_b_tmem.unparse_tokens_mode(tokens, spaced);
307            tokens.push(PtxToken::Comma);
308            if spaced {
309                tokens.push(PtxToken::Space);
310            }
311            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
312            tokens.push(PtxToken::Semicolon);
313            if spaced {
314                tokens.push(PtxToken::Newline);
315            }
316        }
317    }
318
319    impl PtxUnparser for Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsize1 {
320        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
321            self.unparse_tokens_mode(tokens, false);
322        }
323        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
324            push_opcode(tokens, "tcgen05");
325            push_directive(tokens, "mma");
326            match &self.cta_group {
327                CtaGroup::CtaGroup1 => {
328                    push_directive(tokens, "cta_group::1");
329                }
330                CtaGroup::CtaGroup2 => {
331                    push_directive(tokens, "cta_group::2");
332                }
333            }
334            match &self.kind {
335                Kind::KindMxf8f6f4 => {
336                    push_directive(tokens, "kind::mxf8f6f4");
337                }
338                Kind::KindMxf4nvf4 => {
339                    push_directive(tokens, "kind::mxf4nvf4");
340                }
341                Kind::KindMxf4 => {
342                    push_directive(tokens, "kind::mxf4");
343                }
344            }
345            push_directive(tokens, "block_scale");
346            if let Some(scale_vectorsize_5) = self.scale_vectorsize.as_ref() {
347                match scale_vectorsize_5 {
348                    ScaleVectorsize::ScaleVec1x => {
349                        push_directive(tokens, "scale_vec::1X");
350                    }
351                    ScaleVectorsize::ScaleVec2x => {
352                        push_directive(tokens, "scale_vec::2X");
353                    }
354                    ScaleVectorsize::ScaleVec4x => {
355                        push_directive(tokens, "scale_vec::4X");
356                    }
357                    ScaleVectorsize::Block16 => {
358                        push_directive(tokens, "block16");
359                    }
360                    ScaleVectorsize::Block32 => {
361                        push_directive(tokens, "block32");
362                    }
363                }
364            }
365            if spaced {
366                tokens.push(PtxToken::Space);
367            }
368            self.d_tmem.unparse_tokens_mode(tokens, spaced);
369            tokens.push(PtxToken::Comma);
370            if spaced {
371                tokens.push(PtxToken::Space);
372            }
373            self.a_tmem.unparse_tokens_mode(tokens, spaced);
374            tokens.push(PtxToken::Comma);
375            if spaced {
376                tokens.push(PtxToken::Space);
377            }
378            self.b_desc.unparse_tokens_mode(tokens, spaced);
379            tokens.push(PtxToken::Comma);
380            if spaced {
381                tokens.push(PtxToken::Space);
382            }
383            self.idesc.unparse_tokens_mode(tokens, spaced);
384            tokens.push(PtxToken::Comma);
385            if spaced {
386                tokens.push(PtxToken::Space);
387            }
388            self.scale_a_tmem.unparse_tokens_mode(tokens, spaced);
389            tokens.push(PtxToken::Comma);
390            if spaced {
391                tokens.push(PtxToken::Space);
392            }
393            self.scale_b_tmem.unparse_tokens_mode(tokens, spaced);
394            tokens.push(PtxToken::Comma);
395            if spaced {
396                tokens.push(PtxToken::Space);
397            }
398            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
399            tokens.push(PtxToken::Semicolon);
400            if spaced {
401                tokens.push(PtxToken::Newline);
402            }
403        }
404    }
405}
406
407pub mod section_2 {
408    use super::*;
409    use crate::r#type::instruction::tcgen05_mma::section_2::*;
410
411    impl PtxUnparser for Tcgen05MmaCtaGroupKindCollectorUsage {
412        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
413            self.unparse_tokens_mode(tokens, false);
414        }
415        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
416            push_opcode(tokens, "tcgen05");
417            push_directive(tokens, "mma");
418            match &self.cta_group {
419                CtaGroup::CtaGroup1 => {
420                    push_directive(tokens, "cta_group::1");
421                }
422                CtaGroup::CtaGroup2 => {
423                    push_directive(tokens, "cta_group::2");
424                }
425            }
426            match &self.kind {
427                Kind::KindF8f6f4 => {
428                    push_directive(tokens, "kind::f8f6f4");
429                }
430                Kind::KindTf32 => {
431                    push_directive(tokens, "kind::tf32");
432                }
433                Kind::KindF16 => {
434                    push_directive(tokens, "kind::f16");
435                }
436            }
437            match &self.collector_usage {
438                CollectorUsage::CollectorBufferOp(_, n1, n2) => {
439                    let mut combined = String::new();
440                    combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
441                    combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
442                    tokens.push(PtxToken::Dot);
443                    tokens.push(PtxToken::Identifier(
444                        format!("{}{}", "collector", combined).into(),
445                    ));
446                }
447            }
448            if spaced {
449                tokens.push(PtxToken::Space);
450            }
451            self.d_tmem.unparse_tokens_mode(tokens, spaced);
452            tokens.push(PtxToken::Comma);
453            if spaced {
454                tokens.push(PtxToken::Space);
455            }
456            self.a_desc.unparse_tokens_mode(tokens, spaced);
457            tokens.push(PtxToken::Comma);
458            if spaced {
459                tokens.push(PtxToken::Space);
460            }
461            self.b_desc.unparse_tokens_mode(tokens, spaced);
462            tokens.push(PtxToken::Comma);
463            if spaced {
464                tokens.push(PtxToken::Space);
465            }
466            self.idesc.unparse_tokens_mode(tokens, spaced);
467            if self.disable_output_lane.is_some() {
468                tokens.push(PtxToken::Comma);
469            }
470            if let Some(opt_6) = self.disable_output_lane.as_ref() {
471                if spaced {
472                    tokens.push(PtxToken::Space);
473                }
474                opt_6.unparse_tokens_mode(tokens, spaced);
475            }
476            tokens.push(PtxToken::Comma);
477            if spaced {
478                tokens.push(PtxToken::Space);
479            }
480            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
481            if self.scale_input_d.is_some() {
482                tokens.push(PtxToken::Comma);
483            }
484            if let Some(opt_7) = self.scale_input_d.as_ref() {
485                if spaced {
486                    tokens.push(PtxToken::Space);
487                }
488                opt_7.unparse_tokens_mode(tokens, spaced);
489            }
490            tokens.push(PtxToken::Semicolon);
491            if spaced {
492                tokens.push(PtxToken::Newline);
493            }
494        }
495    }
496
497    impl PtxUnparser for Tcgen05MmaCtaGroupKindAshiftCollectorUsage {
498        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
499            self.unparse_tokens_mode(tokens, false);
500        }
501        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
502            push_opcode(tokens, "tcgen05");
503            push_directive(tokens, "mma");
504            match &self.cta_group {
505                CtaGroup::CtaGroup1 => {
506                    push_directive(tokens, "cta_group::1");
507                }
508                CtaGroup::CtaGroup2 => {
509                    push_directive(tokens, "cta_group::2");
510                }
511            }
512            match &self.kind {
513                Kind::KindF8f6f4 => {
514                    push_directive(tokens, "kind::f8f6f4");
515                }
516                Kind::KindTf32 => {
517                    push_directive(tokens, "kind::tf32");
518                }
519                Kind::KindF16 => {
520                    push_directive(tokens, "kind::f16");
521                }
522            }
523            if self.ashift {
524                push_directive(tokens, "ashift");
525            }
526            match &self.collector_usage {
527                CollectorUsage::CollectorBufferOp(_, n1, n2) => {
528                    let mut combined = String::new();
529                    combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
530                    combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
531                    tokens.push(PtxToken::Dot);
532                    tokens.push(PtxToken::Identifier(
533                        format!("{}{}", "collector", combined).into(),
534                    ));
535                }
536            }
537            if spaced {
538                tokens.push(PtxToken::Space);
539            }
540            self.d_tmem.unparse_tokens_mode(tokens, spaced);
541            tokens.push(PtxToken::Comma);
542            if spaced {
543                tokens.push(PtxToken::Space);
544            }
545            self.a_tmem.unparse_tokens_mode(tokens, spaced);
546            tokens.push(PtxToken::Comma);
547            if spaced {
548                tokens.push(PtxToken::Space);
549            }
550            self.b_desc.unparse_tokens_mode(tokens, spaced);
551            tokens.push(PtxToken::Comma);
552            if spaced {
553                tokens.push(PtxToken::Space);
554            }
555            self.idesc.unparse_tokens_mode(tokens, spaced);
556            if self.disable_output_lane.is_some() {
557                tokens.push(PtxToken::Comma);
558            }
559            if let Some(opt_8) = self.disable_output_lane.as_ref() {
560                if spaced {
561                    tokens.push(PtxToken::Space);
562                }
563                opt_8.unparse_tokens_mode(tokens, spaced);
564            }
565            tokens.push(PtxToken::Comma);
566            if spaced {
567                tokens.push(PtxToken::Space);
568            }
569            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
570            if self.scale_input_d.is_some() {
571                tokens.push(PtxToken::Comma);
572            }
573            if let Some(opt_9) = self.scale_input_d.as_ref() {
574                if spaced {
575                    tokens.push(PtxToken::Space);
576                }
577                opt_9.unparse_tokens_mode(tokens, spaced);
578            }
579            tokens.push(PtxToken::Semicolon);
580            if spaced {
581                tokens.push(PtxToken::Newline);
582            }
583        }
584    }
585
586    impl PtxUnparser for Tcgen05MmaCtaGroupKindAshiftCollectorUsage1 {
587        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
588            self.unparse_tokens_mode(tokens, false);
589        }
590        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
591            push_opcode(tokens, "tcgen05");
592            push_directive(tokens, "mma");
593            match &self.cta_group {
594                CtaGroup::CtaGroup1 => {
595                    push_directive(tokens, "cta_group::1");
596                }
597                CtaGroup::CtaGroup2 => {
598                    push_directive(tokens, "cta_group::2");
599                }
600            }
601            match &self.kind {
602                Kind::KindF8f6f4 => {
603                    push_directive(tokens, "kind::f8f6f4");
604                }
605                Kind::KindTf32 => {
606                    push_directive(tokens, "kind::tf32");
607                }
608                Kind::KindF16 => {
609                    push_directive(tokens, "kind::f16");
610                }
611            }
612            push_directive(tokens, "ashift");
613            if let Some(collector_usage_10) = self.collector_usage.as_ref() {
614                match collector_usage_10 {
615                    CollectorUsage::CollectorBufferOp(_, n1, n2) => {
616                        let mut combined = String::new();
617                        combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
618                        combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
619                        tokens.push(PtxToken::Dot);
620                        tokens.push(PtxToken::Identifier(
621                            format!("{}{}", "collector", combined).into(),
622                        ));
623                    }
624                }
625            }
626            if spaced {
627                tokens.push(PtxToken::Space);
628            }
629            self.d_tmem.unparse_tokens_mode(tokens, spaced);
630            tokens.push(PtxToken::Comma);
631            if spaced {
632                tokens.push(PtxToken::Space);
633            }
634            self.a_tmem.unparse_tokens_mode(tokens, spaced);
635            tokens.push(PtxToken::Comma);
636            if spaced {
637                tokens.push(PtxToken::Space);
638            }
639            self.b_desc.unparse_tokens_mode(tokens, spaced);
640            tokens.push(PtxToken::Comma);
641            if spaced {
642                tokens.push(PtxToken::Space);
643            }
644            self.idesc.unparse_tokens_mode(tokens, spaced);
645            if self.disable_output_lane.is_some() {
646                tokens.push(PtxToken::Comma);
647            }
648            if let Some(opt_11) = self.disable_output_lane.as_ref() {
649                if spaced {
650                    tokens.push(PtxToken::Space);
651                }
652                opt_11.unparse_tokens_mode(tokens, spaced);
653            }
654            tokens.push(PtxToken::Comma);
655            if spaced {
656                tokens.push(PtxToken::Space);
657            }
658            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
659            if self.scale_input_d.is_some() {
660                tokens.push(PtxToken::Comma);
661            }
662            if let Some(opt_12) = self.scale_input_d.as_ref() {
663                if spaced {
664                    tokens.push(PtxToken::Space);
665                }
666                opt_12.unparse_tokens_mode(tokens, spaced);
667            }
668            tokens.push(PtxToken::Semicolon);
669            if spaced {
670                tokens.push(PtxToken::Newline);
671            }
672        }
673    }
674}
675
676pub mod section_3 {
677    use super::*;
678    use crate::r#type::instruction::tcgen05_mma::section_3::*;
679
680    impl PtxUnparser for Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage {
681        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
682            self.unparse_tokens_mode(tokens, false);
683        }
684        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
685            push_opcode(tokens, "tcgen05");
686            push_directive(tokens, "mma");
687            match &self.cta_group {
688                CtaGroup::CtaGroup1 => {
689                    push_directive(tokens, "cta_group::1");
690                }
691                CtaGroup::CtaGroup2 => {
692                    push_directive(tokens, "cta_group::2");
693                }
694            }
695            match &self.kind {
696                Kind::KindMxf8f6f4 => {
697                    push_directive(tokens, "kind::mxf8f6f4");
698                }
699                Kind::KindMxf4nvf4 => {
700                    push_directive(tokens, "kind::mxf4nvf4");
701                }
702                Kind::KindMxf4 => {
703                    push_directive(tokens, "kind::mxf4");
704                }
705            }
706            push_directive(tokens, "block_scale");
707            if let Some(scale_vectorsize_13) = self.scale_vectorsize.as_ref() {
708                match scale_vectorsize_13 {
709                    ScaleVectorsize::ScaleVec1x => {
710                        push_directive(tokens, "scale_vec::1X");
711                    }
712                    ScaleVectorsize::ScaleVec2x => {
713                        push_directive(tokens, "scale_vec::2X");
714                    }
715                    ScaleVectorsize::ScaleVec4x => {
716                        push_directive(tokens, "scale_vec::4X");
717                    }
718                    ScaleVectorsize::Block16 => {
719                        push_directive(tokens, "block16");
720                    }
721                    ScaleVectorsize::Block32 => {
722                        push_directive(tokens, "block32");
723                    }
724                }
725            }
726            match &self.collector_usage {
727                CollectorUsage::CollectorBufferOp(_, n1, n2) => {
728                    let mut combined = String::new();
729                    combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
730                    combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
731                    tokens.push(PtxToken::Dot);
732                    tokens.push(PtxToken::Identifier(
733                        format!("{}{}", "collector", combined).into(),
734                    ));
735                }
736            }
737            if spaced {
738                tokens.push(PtxToken::Space);
739            }
740            self.d_tmem.unparse_tokens_mode(tokens, spaced);
741            tokens.push(PtxToken::Comma);
742            if spaced {
743                tokens.push(PtxToken::Space);
744            }
745            self.a_desc.unparse_tokens_mode(tokens, spaced);
746            tokens.push(PtxToken::Comma);
747            if spaced {
748                tokens.push(PtxToken::Space);
749            }
750            self.b_desc.unparse_tokens_mode(tokens, spaced);
751            tokens.push(PtxToken::Comma);
752            if spaced {
753                tokens.push(PtxToken::Space);
754            }
755            self.idesc.unparse_tokens_mode(tokens, spaced);
756            tokens.push(PtxToken::Comma);
757            if spaced {
758                tokens.push(PtxToken::Space);
759            }
760            self.scale_a_tmem.unparse_tokens_mode(tokens, spaced);
761            tokens.push(PtxToken::Comma);
762            if spaced {
763                tokens.push(PtxToken::Space);
764            }
765            self.scale_b_tmem.unparse_tokens_mode(tokens, spaced);
766            tokens.push(PtxToken::Comma);
767            if spaced {
768                tokens.push(PtxToken::Space);
769            }
770            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
771            tokens.push(PtxToken::Semicolon);
772            if spaced {
773                tokens.push(PtxToken::Newline);
774            }
775        }
776    }
777
778    impl PtxUnparser for Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage1 {
779        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
780            self.unparse_tokens_mode(tokens, false);
781        }
782        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
783            push_opcode(tokens, "tcgen05");
784            push_directive(tokens, "mma");
785            match &self.cta_group {
786                CtaGroup::CtaGroup1 => {
787                    push_directive(tokens, "cta_group::1");
788                }
789                CtaGroup::CtaGroup2 => {
790                    push_directive(tokens, "cta_group::2");
791                }
792            }
793            match &self.kind {
794                Kind::KindMxf8f6f4 => {
795                    push_directive(tokens, "kind::mxf8f6f4");
796                }
797                Kind::KindMxf4nvf4 => {
798                    push_directive(tokens, "kind::mxf4nvf4");
799                }
800                Kind::KindMxf4 => {
801                    push_directive(tokens, "kind::mxf4");
802                }
803            }
804            push_directive(tokens, "block_scale");
805            if let Some(scale_vectorsize_14) = self.scale_vectorsize.as_ref() {
806                match scale_vectorsize_14 {
807                    ScaleVectorsize::ScaleVec1x => {
808                        push_directive(tokens, "scale_vec::1X");
809                    }
810                    ScaleVectorsize::ScaleVec2x => {
811                        push_directive(tokens, "scale_vec::2X");
812                    }
813                    ScaleVectorsize::ScaleVec4x => {
814                        push_directive(tokens, "scale_vec::4X");
815                    }
816                    ScaleVectorsize::Block16 => {
817                        push_directive(tokens, "block16");
818                    }
819                    ScaleVectorsize::Block32 => {
820                        push_directive(tokens, "block32");
821                    }
822                }
823            }
824            match &self.collector_usage {
825                CollectorUsage::CollectorBufferOp(_, n1, n2) => {
826                    let mut combined = String::new();
827                    combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
828                    combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
829                    tokens.push(PtxToken::Dot);
830                    tokens.push(PtxToken::Identifier(
831                        format!("{}{}", "collector", combined).into(),
832                    ));
833                }
834            }
835            if spaced {
836                tokens.push(PtxToken::Space);
837            }
838            self.d_tmem.unparse_tokens_mode(tokens, spaced);
839            tokens.push(PtxToken::Comma);
840            if spaced {
841                tokens.push(PtxToken::Space);
842            }
843            self.a_tmem.unparse_tokens_mode(tokens, spaced);
844            tokens.push(PtxToken::Comma);
845            if spaced {
846                tokens.push(PtxToken::Space);
847            }
848            self.b_desc.unparse_tokens_mode(tokens, spaced);
849            tokens.push(PtxToken::Comma);
850            if spaced {
851                tokens.push(PtxToken::Space);
852            }
853            self.idesc.unparse_tokens_mode(tokens, spaced);
854            tokens.push(PtxToken::Comma);
855            if spaced {
856                tokens.push(PtxToken::Space);
857            }
858            self.scale_a_tmem.unparse_tokens_mode(tokens, spaced);
859            tokens.push(PtxToken::Comma);
860            if spaced {
861                tokens.push(PtxToken::Space);
862            }
863            self.scale_b_tmem.unparse_tokens_mode(tokens, spaced);
864            tokens.push(PtxToken::Comma);
865            if spaced {
866                tokens.push(PtxToken::Space);
867            }
868            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
869            tokens.push(PtxToken::Semicolon);
870            if spaced {
871                tokens.push(PtxToken::Newline);
872            }
873        }
874    }
875}
876
877pub mod section_4 {
878    use super::*;
879    use crate::r#type::instruction::tcgen05_mma::section_4::*;
880
881    impl PtxUnparser for Tcgen05MmaCtaGroupKindI8 {
882        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
883            self.unparse_tokens_mode(tokens, false);
884        }
885        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
886            push_opcode(tokens, "tcgen05");
887            push_directive(tokens, "mma");
888            match &self.cta_group {
889                CtaGroup::CtaGroup1 => {
890                    push_directive(tokens, "cta_group::1");
891                }
892                CtaGroup::CtaGroup2 => {
893                    push_directive(tokens, "cta_group::2");
894                }
895            }
896            push_directive(tokens, "kind::i8");
897            if spaced {
898                tokens.push(PtxToken::Space);
899            }
900            self.d_tmem.unparse_tokens_mode(tokens, spaced);
901            tokens.push(PtxToken::Comma);
902            if spaced {
903                tokens.push(PtxToken::Space);
904            }
905            self.a_desc.unparse_tokens_mode(tokens, spaced);
906            tokens.push(PtxToken::Comma);
907            if spaced {
908                tokens.push(PtxToken::Space);
909            }
910            self.b_desc.unparse_tokens_mode(tokens, spaced);
911            tokens.push(PtxToken::Comma);
912            if spaced {
913                tokens.push(PtxToken::Space);
914            }
915            self.idesc.unparse_tokens_mode(tokens, spaced);
916            if self.disable_output_lane.is_some() {
917                tokens.push(PtxToken::Comma);
918            }
919            if let Some(opt_15) = self.disable_output_lane.as_ref() {
920                if spaced {
921                    tokens.push(PtxToken::Space);
922                }
923                opt_15.unparse_tokens_mode(tokens, spaced);
924            }
925            tokens.push(PtxToken::Comma);
926            if spaced {
927                tokens.push(PtxToken::Space);
928            }
929            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
930            tokens.push(PtxToken::Semicolon);
931            if spaced {
932                tokens.push(PtxToken::Newline);
933            }
934        }
935    }
936
937    impl PtxUnparser for Tcgen05MmaCtaGroupKindI81 {
938        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
939            self.unparse_tokens_mode(tokens, false);
940        }
941        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
942            push_opcode(tokens, "tcgen05");
943            push_directive(tokens, "mma");
944            match &self.cta_group {
945                CtaGroup::CtaGroup1 => {
946                    push_directive(tokens, "cta_group::1");
947                }
948                CtaGroup::CtaGroup2 => {
949                    push_directive(tokens, "cta_group::2");
950                }
951            }
952            push_directive(tokens, "kind::i8");
953            if spaced {
954                tokens.push(PtxToken::Space);
955            }
956            self.d_tmem.unparse_tokens_mode(tokens, spaced);
957            tokens.push(PtxToken::Comma);
958            if spaced {
959                tokens.push(PtxToken::Space);
960            }
961            self.a_tmem.unparse_tokens_mode(tokens, spaced);
962            tokens.push(PtxToken::Comma);
963            if spaced {
964                tokens.push(PtxToken::Space);
965            }
966            self.b_desc.unparse_tokens_mode(tokens, spaced);
967            tokens.push(PtxToken::Comma);
968            if spaced {
969                tokens.push(PtxToken::Space);
970            }
971            self.idesc.unparse_tokens_mode(tokens, spaced);
972            if self.disable_output_lane.is_some() {
973                tokens.push(PtxToken::Comma);
974            }
975            if let Some(opt_16) = self.disable_output_lane.as_ref() {
976                if spaced {
977                    tokens.push(PtxToken::Space);
978                }
979                opt_16.unparse_tokens_mode(tokens, spaced);
980            }
981            tokens.push(PtxToken::Comma);
982            if spaced {
983                tokens.push(PtxToken::Space);
984            }
985            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
986            tokens.push(PtxToken::Semicolon);
987            if spaced {
988                tokens.push(PtxToken::Newline);
989            }
990        }
991    }
992}
993
994pub mod section_5 {
995    use super::*;
996    use crate::r#type::instruction::tcgen05_mma::section_5::*;
997
998    impl PtxUnparser for Tcgen05MmaCtaGroupKindI8CollectorUsage {
999        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
1000            self.unparse_tokens_mode(tokens, false);
1001        }
1002        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
1003            push_opcode(tokens, "tcgen05");
1004            push_directive(tokens, "mma");
1005            match &self.cta_group {
1006                CtaGroup::CtaGroup1 => {
1007                    push_directive(tokens, "cta_group::1");
1008                }
1009                CtaGroup::CtaGroup2 => {
1010                    push_directive(tokens, "cta_group::2");
1011                }
1012            }
1013            push_directive(tokens, "kind::i8");
1014            match &self.collector_usage {
1015                CollectorUsage::CollectorBufferOp(_, n1, n2) => {
1016                    let mut combined = String::new();
1017                    combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
1018                    combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
1019                    tokens.push(PtxToken::Dot);
1020                    tokens.push(PtxToken::Identifier(
1021                        format!("{}{}", "collector", combined).into(),
1022                    ));
1023                }
1024            }
1025            if spaced {
1026                tokens.push(PtxToken::Space);
1027            }
1028            self.d_tmem.unparse_tokens_mode(tokens, spaced);
1029            tokens.push(PtxToken::Comma);
1030            if spaced {
1031                tokens.push(PtxToken::Space);
1032            }
1033            self.a_desc.unparse_tokens_mode(tokens, spaced);
1034            tokens.push(PtxToken::Comma);
1035            if spaced {
1036                tokens.push(PtxToken::Space);
1037            }
1038            self.b_desc.unparse_tokens_mode(tokens, spaced);
1039            tokens.push(PtxToken::Comma);
1040            if spaced {
1041                tokens.push(PtxToken::Space);
1042            }
1043            self.idesc.unparse_tokens_mode(tokens, spaced);
1044            if self.disable_output_lane.is_some() {
1045                tokens.push(PtxToken::Comma);
1046            }
1047            if let Some(opt_17) = self.disable_output_lane.as_ref() {
1048                if spaced {
1049                    tokens.push(PtxToken::Space);
1050                }
1051                opt_17.unparse_tokens_mode(tokens, spaced);
1052            }
1053            tokens.push(PtxToken::Comma);
1054            if spaced {
1055                tokens.push(PtxToken::Space);
1056            }
1057            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
1058            tokens.push(PtxToken::Semicolon);
1059            if spaced {
1060                tokens.push(PtxToken::Newline);
1061            }
1062        }
1063    }
1064
1065    impl PtxUnparser for Tcgen05MmaCtaGroupKindI8AshiftCollectorUsage {
1066        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
1067            self.unparse_tokens_mode(tokens, false);
1068        }
1069        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
1070            push_opcode(tokens, "tcgen05");
1071            push_directive(tokens, "mma");
1072            match &self.cta_group {
1073                CtaGroup::CtaGroup1 => {
1074                    push_directive(tokens, "cta_group::1");
1075                }
1076                CtaGroup::CtaGroup2 => {
1077                    push_directive(tokens, "cta_group::2");
1078                }
1079            }
1080            push_directive(tokens, "kind::i8");
1081            push_directive(tokens, "ashift");
1082            if let Some(collector_usage_18) = self.collector_usage.as_ref() {
1083                match collector_usage_18 {
1084                    CollectorUsage::CollectorBufferOp(_, n1, n2) => {
1085                        let mut combined = String::new();
1086                        combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
1087                        combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
1088                        tokens.push(PtxToken::Dot);
1089                        tokens.push(PtxToken::Identifier(
1090                            format!("{}{}", "collector", combined).into(),
1091                        ));
1092                    }
1093                }
1094            }
1095            if spaced {
1096                tokens.push(PtxToken::Space);
1097            }
1098            self.d_tmem.unparse_tokens_mode(tokens, spaced);
1099            tokens.push(PtxToken::Comma);
1100            if spaced {
1101                tokens.push(PtxToken::Space);
1102            }
1103            self.a_tmem.unparse_tokens_mode(tokens, spaced);
1104            tokens.push(PtxToken::Comma);
1105            if spaced {
1106                tokens.push(PtxToken::Space);
1107            }
1108            self.b_desc.unparse_tokens_mode(tokens, spaced);
1109            tokens.push(PtxToken::Comma);
1110            if spaced {
1111                tokens.push(PtxToken::Space);
1112            }
1113            self.idesc.unparse_tokens_mode(tokens, spaced);
1114            if self.disable_output_lane.is_some() {
1115                tokens.push(PtxToken::Comma);
1116            }
1117            if let Some(opt_19) = self.disable_output_lane.as_ref() {
1118                if spaced {
1119                    tokens.push(PtxToken::Space);
1120                }
1121                opt_19.unparse_tokens_mode(tokens, spaced);
1122            }
1123            tokens.push(PtxToken::Comma);
1124            if spaced {
1125                tokens.push(PtxToken::Space);
1126            }
1127            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
1128            tokens.push(PtxToken::Semicolon);
1129            if spaced {
1130                tokens.push(PtxToken::Newline);
1131            }
1132        }
1133    }
1134
1135    impl PtxUnparser for Tcgen05MmaCtaGroupKindI8AshiftCollectorUsage1 {
1136        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
1137            self.unparse_tokens_mode(tokens, false);
1138        }
1139        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
1140            push_opcode(tokens, "tcgen05");
1141            push_directive(tokens, "mma");
1142            match &self.cta_group {
1143                CtaGroup::CtaGroup1 => {
1144                    push_directive(tokens, "cta_group::1");
1145                }
1146                CtaGroup::CtaGroup2 => {
1147                    push_directive(tokens, "cta_group::2");
1148                }
1149            }
1150            push_directive(tokens, "kind::i8");
1151            if self.ashift {
1152                push_directive(tokens, "ashift");
1153            }
1154            match &self.collector_usage {
1155                CollectorUsage::CollectorBufferOp(_, n1, n2) => {
1156                    let mut combined = String::new();
1157                    combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
1158                    combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
1159                    tokens.push(PtxToken::Dot);
1160                    tokens.push(PtxToken::Identifier(
1161                        format!("{}{}", "collector", combined).into(),
1162                    ));
1163                }
1164            }
1165            if spaced {
1166                tokens.push(PtxToken::Space);
1167            }
1168            self.d_tmem.unparse_tokens_mode(tokens, spaced);
1169            tokens.push(PtxToken::Comma);
1170            if spaced {
1171                tokens.push(PtxToken::Space);
1172            }
1173            self.a_tmem.unparse_tokens_mode(tokens, spaced);
1174            tokens.push(PtxToken::Comma);
1175            if spaced {
1176                tokens.push(PtxToken::Space);
1177            }
1178            self.b_desc.unparse_tokens_mode(tokens, spaced);
1179            tokens.push(PtxToken::Comma);
1180            if spaced {
1181                tokens.push(PtxToken::Space);
1182            }
1183            self.idesc.unparse_tokens_mode(tokens, spaced);
1184            if self.disable_output_lane.is_some() {
1185                tokens.push(PtxToken::Comma);
1186            }
1187            if let Some(opt_20) = self.disable_output_lane.as_ref() {
1188                if spaced {
1189                    tokens.push(PtxToken::Space);
1190                }
1191                opt_20.unparse_tokens_mode(tokens, spaced);
1192            }
1193            tokens.push(PtxToken::Comma);
1194            if spaced {
1195                tokens.push(PtxToken::Space);
1196            }
1197            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
1198            tokens.push(PtxToken::Semicolon);
1199            if spaced {
1200                tokens.push(PtxToken::Newline);
1201            }
1202        }
1203    }
1204}
ptx_parser/unparser/instruction/tcgen05_mma.rs

ptx_parser/unparser/instruction/
tcgen05_mma.rs
