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            push_opcode(tokens, "tcgen05");
80            push_directive(tokens, "mma");
81            match &self.cta_group {
82                CtaGroup::CtaGroup1 => {
83                    push_directive(tokens, "cta_group::1");
84                }
85                CtaGroup::CtaGroup2 => {
86                    push_directive(tokens, "cta_group::2");
87                }
88            }
89            match &self.kind {
90                Kind::KindF8f6f4 => {
91                    push_directive(tokens, "kind::f8f6f4");
92                }
93                Kind::KindTf32 => {
94                    push_directive(tokens, "kind::tf32");
95                }
96                Kind::KindF16 => {
97                    push_directive(tokens, "kind::f16");
98                }
99            }
100            self.d_tmem.unparse_tokens(tokens);
101            tokens.push(PtxToken::Comma);
102            self.a_desc.unparse_tokens(tokens);
103            tokens.push(PtxToken::Comma);
104            self.b_desc.unparse_tokens(tokens);
105            tokens.push(PtxToken::Comma);
106            self.idesc.unparse_tokens(tokens);
107            if self.disable_output_lane.is_some() {
108                tokens.push(PtxToken::Comma);
109            }
110            if let Some(opt_0) = self.disable_output_lane.as_ref() {
111                opt_0.unparse_tokens(tokens);
112            }
113            tokens.push(PtxToken::Comma);
114            self.enable_input_d.unparse_tokens(tokens);
115            if self.scale_input_d.is_some() {
116                tokens.push(PtxToken::Comma);
117            }
118            if let Some(opt_1) = self.scale_input_d.as_ref() {
119                opt_1.unparse_tokens(tokens);
120            }
121            tokens.push(PtxToken::Semicolon);
122        }
123    }
124
125    impl PtxUnparser for Tcgen05MmaCtaGroupKind1 {
126        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
127            push_opcode(tokens, "tcgen05");
128            push_directive(tokens, "mma");
129            match &self.cta_group {
130                CtaGroup::CtaGroup1 => {
131                    push_directive(tokens, "cta_group::1");
132                }
133                CtaGroup::CtaGroup2 => {
134                    push_directive(tokens, "cta_group::2");
135                }
136            }
137            match &self.kind {
138                Kind::KindF8f6f4 => {
139                    push_directive(tokens, "kind::f8f6f4");
140                }
141                Kind::KindTf32 => {
142                    push_directive(tokens, "kind::tf32");
143                }
144                Kind::KindF16 => {
145                    push_directive(tokens, "kind::f16");
146                }
147            }
148            self.d_tmem.unparse_tokens(tokens);
149            tokens.push(PtxToken::Comma);
150            self.a_tmem.unparse_tokens(tokens);
151            tokens.push(PtxToken::Comma);
152            self.b_desc.unparse_tokens(tokens);
153            tokens.push(PtxToken::Comma);
154            self.idesc.unparse_tokens(tokens);
155            if self.disable_output_lane.is_some() {
156                tokens.push(PtxToken::Comma);
157            }
158            if let Some(opt_2) = self.disable_output_lane.as_ref() {
159                opt_2.unparse_tokens(tokens);
160            }
161            tokens.push(PtxToken::Comma);
162            self.enable_input_d.unparse_tokens(tokens);
163            if self.scale_input_d.is_some() {
164                tokens.push(PtxToken::Comma);
165            }
166            if let Some(opt_3) = self.scale_input_d.as_ref() {
167                opt_3.unparse_tokens(tokens);
168            }
169            tokens.push(PtxToken::Semicolon);
170        }
171    }
172}
173
174pub mod section_1 {
175    use super::*;
176    use crate::r#type::instruction::tcgen05_mma::section_1::*;
177
178    impl PtxUnparser for Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsize {
179        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
180            push_opcode(tokens, "tcgen05");
181            push_directive(tokens, "mma");
182            match &self.cta_group {
183                CtaGroup::CtaGroup1 => {
184                    push_directive(tokens, "cta_group::1");
185                }
186                CtaGroup::CtaGroup2 => {
187                    push_directive(tokens, "cta_group::2");
188                }
189            }
190            match &self.kind {
191                Kind::KindMxf8f6f4 => {
192                    push_directive(tokens, "kind::mxf8f6f4");
193                }
194                Kind::KindMxf4nvf4 => {
195                    push_directive(tokens, "kind::mxf4nvf4");
196                }
197                Kind::KindMxf4 => {
198                    push_directive(tokens, "kind::mxf4");
199                }
200            }
201            push_directive(tokens, "block_scale");
202            if let Some(scale_vectorsize_4) = self.scale_vectorsize.as_ref() {
203                match scale_vectorsize_4 {
204                    ScaleVectorsize::ScaleVec1x => {
205                        push_directive(tokens, "scale_vec::1X");
206                    }
207                    ScaleVectorsize::ScaleVec2x => {
208                        push_directive(tokens, "scale_vec::2X");
209                    }
210                    ScaleVectorsize::ScaleVec4x => {
211                        push_directive(tokens, "scale_vec::4X");
212                    }
213                    ScaleVectorsize::Block16 => {
214                        push_directive(tokens, "block16");
215                    }
216                    ScaleVectorsize::Block32 => {
217                        push_directive(tokens, "block32");
218                    }
219                }
220            }
221            self.d_tmem.unparse_tokens(tokens);
222            tokens.push(PtxToken::Comma);
223            self.a_desc.unparse_tokens(tokens);
224            tokens.push(PtxToken::Comma);
225            self.b_desc.unparse_tokens(tokens);
226            tokens.push(PtxToken::Comma);
227            self.idesc.unparse_tokens(tokens);
228            tokens.push(PtxToken::Comma);
229            self.scale_a_tmem.unparse_tokens(tokens);
230            tokens.push(PtxToken::Comma);
231            self.scale_b_tmem.unparse_tokens(tokens);
232            tokens.push(PtxToken::Comma);
233            self.enable_input_d.unparse_tokens(tokens);
234            tokens.push(PtxToken::Semicolon);
235        }
236    }
237
238    impl PtxUnparser for Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsize1 {
239        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
240            push_opcode(tokens, "tcgen05");
241            push_directive(tokens, "mma");
242            match &self.cta_group {
243                CtaGroup::CtaGroup1 => {
244                    push_directive(tokens, "cta_group::1");
245                }
246                CtaGroup::CtaGroup2 => {
247                    push_directive(tokens, "cta_group::2");
248                }
249            }
250            match &self.kind {
251                Kind::KindMxf8f6f4 => {
252                    push_directive(tokens, "kind::mxf8f6f4");
253                }
254                Kind::KindMxf4nvf4 => {
255                    push_directive(tokens, "kind::mxf4nvf4");
256                }
257                Kind::KindMxf4 => {
258                    push_directive(tokens, "kind::mxf4");
259                }
260            }
261            push_directive(tokens, "block_scale");
262            if let Some(scale_vectorsize_5) = self.scale_vectorsize.as_ref() {
263                match scale_vectorsize_5 {
264                    ScaleVectorsize::ScaleVec1x => {
265                        push_directive(tokens, "scale_vec::1X");
266                    }
267                    ScaleVectorsize::ScaleVec2x => {
268                        push_directive(tokens, "scale_vec::2X");
269                    }
270                    ScaleVectorsize::ScaleVec4x => {
271                        push_directive(tokens, "scale_vec::4X");
272                    }
273                    ScaleVectorsize::Block16 => {
274                        push_directive(tokens, "block16");
275                    }
276                    ScaleVectorsize::Block32 => {
277                        push_directive(tokens, "block32");
278                    }
279                }
280            }
281            self.d_tmem.unparse_tokens(tokens);
282            tokens.push(PtxToken::Comma);
283            self.a_tmem.unparse_tokens(tokens);
284            tokens.push(PtxToken::Comma);
285            self.b_desc.unparse_tokens(tokens);
286            tokens.push(PtxToken::Comma);
287            self.idesc.unparse_tokens(tokens);
288            tokens.push(PtxToken::Comma);
289            self.scale_a_tmem.unparse_tokens(tokens);
290            tokens.push(PtxToken::Comma);
291            self.scale_b_tmem.unparse_tokens(tokens);
292            tokens.push(PtxToken::Comma);
293            self.enable_input_d.unparse_tokens(tokens);
294            tokens.push(PtxToken::Semicolon);
295        }
296    }
297}
298
299pub mod section_2 {
300    use super::*;
301    use crate::r#type::instruction::tcgen05_mma::section_2::*;
302
303    impl PtxUnparser for Tcgen05MmaCtaGroupKindCollectorUsage {
304        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
305            push_opcode(tokens, "tcgen05");
306            push_directive(tokens, "mma");
307            match &self.cta_group {
308                CtaGroup::CtaGroup1 => {
309                    push_directive(tokens, "cta_group::1");
310                }
311                CtaGroup::CtaGroup2 => {
312                    push_directive(tokens, "cta_group::2");
313                }
314            }
315            match &self.kind {
316                Kind::KindF8f6f4 => {
317                    push_directive(tokens, "kind::f8f6f4");
318                }
319                Kind::KindTf32 => {
320                    push_directive(tokens, "kind::tf32");
321                }
322                Kind::KindF16 => {
323                    push_directive(tokens, "kind::f16");
324                }
325            }
326            match &self.collector_usage {
327                CollectorUsage::CollectorBufferOp(_, n1, n2) => {
328                    let mut combined = String::new();
329                    combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
330                    combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
331                    tokens.push(PtxToken::Dot);
332                    tokens.push(PtxToken::Identifier(
333                        format!("{}{}", "collector", combined).into(),
334                    ));
335                }
336            }
337            self.d_tmem.unparse_tokens(tokens);
338            tokens.push(PtxToken::Comma);
339            self.a_desc.unparse_tokens(tokens);
340            tokens.push(PtxToken::Comma);
341            self.b_desc.unparse_tokens(tokens);
342            tokens.push(PtxToken::Comma);
343            self.idesc.unparse_tokens(tokens);
344            if self.disable_output_lane.is_some() {
345                tokens.push(PtxToken::Comma);
346            }
347            if let Some(opt_6) = self.disable_output_lane.as_ref() {
348                opt_6.unparse_tokens(tokens);
349            }
350            tokens.push(PtxToken::Comma);
351            self.enable_input_d.unparse_tokens(tokens);
352            if self.scale_input_d.is_some() {
353                tokens.push(PtxToken::Comma);
354            }
355            if let Some(opt_7) = self.scale_input_d.as_ref() {
356                opt_7.unparse_tokens(tokens);
357            }
358            tokens.push(PtxToken::Semicolon);
359        }
360    }
361
362    impl PtxUnparser for Tcgen05MmaCtaGroupKindAshiftCollectorUsage {
363        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
364            push_opcode(tokens, "tcgen05");
365            push_directive(tokens, "mma");
366            match &self.cta_group {
367                CtaGroup::CtaGroup1 => {
368                    push_directive(tokens, "cta_group::1");
369                }
370                CtaGroup::CtaGroup2 => {
371                    push_directive(tokens, "cta_group::2");
372                }
373            }
374            match &self.kind {
375                Kind::KindF8f6f4 => {
376                    push_directive(tokens, "kind::f8f6f4");
377                }
378                Kind::KindTf32 => {
379                    push_directive(tokens, "kind::tf32");
380                }
381                Kind::KindF16 => {
382                    push_directive(tokens, "kind::f16");
383                }
384            }
385            if self.ashift {
386                push_directive(tokens, "ashift");
387            }
388            match &self.collector_usage {
389                CollectorUsage::CollectorBufferOp(_, n1, n2) => {
390                    let mut combined = String::new();
391                    combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
392                    combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
393                    tokens.push(PtxToken::Dot);
394                    tokens.push(PtxToken::Identifier(
395                        format!("{}{}", "collector", combined).into(),
396                    ));
397                }
398            }
399            self.d_tmem.unparse_tokens(tokens);
400            tokens.push(PtxToken::Comma);
401            self.a_tmem.unparse_tokens(tokens);
402            tokens.push(PtxToken::Comma);
403            self.b_desc.unparse_tokens(tokens);
404            tokens.push(PtxToken::Comma);
405            self.idesc.unparse_tokens(tokens);
406            if self.disable_output_lane.is_some() {
407                tokens.push(PtxToken::Comma);
408            }
409            if let Some(opt_8) = self.disable_output_lane.as_ref() {
410                opt_8.unparse_tokens(tokens);
411            }
412            tokens.push(PtxToken::Comma);
413            self.enable_input_d.unparse_tokens(tokens);
414            if self.scale_input_d.is_some() {
415                tokens.push(PtxToken::Comma);
416            }
417            if let Some(opt_9) = self.scale_input_d.as_ref() {
418                opt_9.unparse_tokens(tokens);
419            }
420            tokens.push(PtxToken::Semicolon);
421        }
422    }
423
424    impl PtxUnparser for Tcgen05MmaCtaGroupKindAshiftCollectorUsage1 {
425        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
426            push_opcode(tokens, "tcgen05");
427            push_directive(tokens, "mma");
428            match &self.cta_group {
429                CtaGroup::CtaGroup1 => {
430                    push_directive(tokens, "cta_group::1");
431                }
432                CtaGroup::CtaGroup2 => {
433                    push_directive(tokens, "cta_group::2");
434                }
435            }
436            match &self.kind {
437                Kind::KindF8f6f4 => {
438                    push_directive(tokens, "kind::f8f6f4");
439                }
440                Kind::KindTf32 => {
441                    push_directive(tokens, "kind::tf32");
442                }
443                Kind::KindF16 => {
444                    push_directive(tokens, "kind::f16");
445                }
446            }
447            push_directive(tokens, "ashift");
448            if let Some(collector_usage_10) = self.collector_usage.as_ref() {
449                match collector_usage_10 {
450                    CollectorUsage::CollectorBufferOp(_, n1, n2) => {
451                        let mut combined = String::new();
452                        combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
453                        combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
454                        tokens.push(PtxToken::Dot);
455                        tokens.push(PtxToken::Identifier(
456                            format!("{}{}", "collector", combined).into(),
457                        ));
458                    }
459                }
460            }
461            self.d_tmem.unparse_tokens(tokens);
462            tokens.push(PtxToken::Comma);
463            self.a_tmem.unparse_tokens(tokens);
464            tokens.push(PtxToken::Comma);
465            self.b_desc.unparse_tokens(tokens);
466            tokens.push(PtxToken::Comma);
467            self.idesc.unparse_tokens(tokens);
468            if self.disable_output_lane.is_some() {
469                tokens.push(PtxToken::Comma);
470            }
471            if let Some(opt_11) = self.disable_output_lane.as_ref() {
472                opt_11.unparse_tokens(tokens);
473            }
474            tokens.push(PtxToken::Comma);
475            self.enable_input_d.unparse_tokens(tokens);
476            if self.scale_input_d.is_some() {
477                tokens.push(PtxToken::Comma);
478            }
479            if let Some(opt_12) = self.scale_input_d.as_ref() {
480                opt_12.unparse_tokens(tokens);
481            }
482            tokens.push(PtxToken::Semicolon);
483        }
484    }
485}
486
487pub mod section_3 {
488    use super::*;
489    use crate::r#type::instruction::tcgen05_mma::section_3::*;
490
491    impl PtxUnparser for Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage {
492        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
493            push_opcode(tokens, "tcgen05");
494            push_directive(tokens, "mma");
495            match &self.cta_group {
496                CtaGroup::CtaGroup1 => {
497                    push_directive(tokens, "cta_group::1");
498                }
499                CtaGroup::CtaGroup2 => {
500                    push_directive(tokens, "cta_group::2");
501                }
502            }
503            match &self.kind {
504                Kind::KindMxf8f6f4 => {
505                    push_directive(tokens, "kind::mxf8f6f4");
506                }
507                Kind::KindMxf4nvf4 => {
508                    push_directive(tokens, "kind::mxf4nvf4");
509                }
510                Kind::KindMxf4 => {
511                    push_directive(tokens, "kind::mxf4");
512                }
513            }
514            push_directive(tokens, "block_scale");
515            if let Some(scale_vectorsize_13) = self.scale_vectorsize.as_ref() {
516                match scale_vectorsize_13 {
517                    ScaleVectorsize::ScaleVec1x => {
518                        push_directive(tokens, "scale_vec::1X");
519                    }
520                    ScaleVectorsize::ScaleVec2x => {
521                        push_directive(tokens, "scale_vec::2X");
522                    }
523                    ScaleVectorsize::ScaleVec4x => {
524                        push_directive(tokens, "scale_vec::4X");
525                    }
526                    ScaleVectorsize::Block16 => {
527                        push_directive(tokens, "block16");
528                    }
529                    ScaleVectorsize::Block32 => {
530                        push_directive(tokens, "block32");
531                    }
532                }
533            }
534            match &self.collector_usage {
535                CollectorUsage::CollectorBufferOp(_, n1, n2) => {
536                    let mut combined = String::new();
537                    combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
538                    combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
539                    tokens.push(PtxToken::Dot);
540                    tokens.push(PtxToken::Identifier(
541                        format!("{}{}", "collector", combined).into(),
542                    ));
543                }
544            }
545            self.d_tmem.unparse_tokens(tokens);
546            tokens.push(PtxToken::Comma);
547            self.a_desc.unparse_tokens(tokens);
548            tokens.push(PtxToken::Comma);
549            self.b_desc.unparse_tokens(tokens);
550            tokens.push(PtxToken::Comma);
551            self.idesc.unparse_tokens(tokens);
552            tokens.push(PtxToken::Comma);
553            self.scale_a_tmem.unparse_tokens(tokens);
554            tokens.push(PtxToken::Comma);
555            self.scale_b_tmem.unparse_tokens(tokens);
556            tokens.push(PtxToken::Comma);
557            self.enable_input_d.unparse_tokens(tokens);
558            tokens.push(PtxToken::Semicolon);
559        }
560    }
561
562    impl PtxUnparser for Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage1 {
563        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
564            push_opcode(tokens, "tcgen05");
565            push_directive(tokens, "mma");
566            match &self.cta_group {
567                CtaGroup::CtaGroup1 => {
568                    push_directive(tokens, "cta_group::1");
569                }
570                CtaGroup::CtaGroup2 => {
571                    push_directive(tokens, "cta_group::2");
572                }
573            }
574            match &self.kind {
575                Kind::KindMxf8f6f4 => {
576                    push_directive(tokens, "kind::mxf8f6f4");
577                }
578                Kind::KindMxf4nvf4 => {
579                    push_directive(tokens, "kind::mxf4nvf4");
580                }
581                Kind::KindMxf4 => {
582                    push_directive(tokens, "kind::mxf4");
583                }
584            }
585            push_directive(tokens, "block_scale");
586            if let Some(scale_vectorsize_14) = self.scale_vectorsize.as_ref() {
587                match scale_vectorsize_14 {
588                    ScaleVectorsize::ScaleVec1x => {
589                        push_directive(tokens, "scale_vec::1X");
590                    }
591                    ScaleVectorsize::ScaleVec2x => {
592                        push_directive(tokens, "scale_vec::2X");
593                    }
594                    ScaleVectorsize::ScaleVec4x => {
595                        push_directive(tokens, "scale_vec::4X");
596                    }
597                    ScaleVectorsize::Block16 => {
598                        push_directive(tokens, "block16");
599                    }
600                    ScaleVectorsize::Block32 => {
601                        push_directive(tokens, "block32");
602                    }
603                }
604            }
605            match &self.collector_usage {
606                CollectorUsage::CollectorBufferOp(_, n1, n2) => {
607                    let mut combined = String::new();
608                    combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
609                    combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
610                    tokens.push(PtxToken::Dot);
611                    tokens.push(PtxToken::Identifier(
612                        format!("{}{}", "collector", combined).into(),
613                    ));
614                }
615            }
616            self.d_tmem.unparse_tokens(tokens);
617            tokens.push(PtxToken::Comma);
618            self.a_tmem.unparse_tokens(tokens);
619            tokens.push(PtxToken::Comma);
620            self.b_desc.unparse_tokens(tokens);
621            tokens.push(PtxToken::Comma);
622            self.idesc.unparse_tokens(tokens);
623            tokens.push(PtxToken::Comma);
624            self.scale_a_tmem.unparse_tokens(tokens);
625            tokens.push(PtxToken::Comma);
626            self.scale_b_tmem.unparse_tokens(tokens);
627            tokens.push(PtxToken::Comma);
628            self.enable_input_d.unparse_tokens(tokens);
629            tokens.push(PtxToken::Semicolon);
630        }
631    }
632}
633
634pub mod section_4 {
635    use super::*;
636    use crate::r#type::instruction::tcgen05_mma::section_4::*;
637
638    impl PtxUnparser for Tcgen05MmaCtaGroupKindI8 {
639        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
640            push_opcode(tokens, "tcgen05");
641            push_directive(tokens, "mma");
642            match &self.cta_group {
643                CtaGroup::CtaGroup1 => {
644                    push_directive(tokens, "cta_group::1");
645                }
646                CtaGroup::CtaGroup2 => {
647                    push_directive(tokens, "cta_group::2");
648                }
649            }
650            push_directive(tokens, "kind::i8");
651            self.d_tmem.unparse_tokens(tokens);
652            tokens.push(PtxToken::Comma);
653            self.a_desc.unparse_tokens(tokens);
654            tokens.push(PtxToken::Comma);
655            self.b_desc.unparse_tokens(tokens);
656            tokens.push(PtxToken::Comma);
657            self.idesc.unparse_tokens(tokens);
658            if self.disable_output_lane.is_some() {
659                tokens.push(PtxToken::Comma);
660            }
661            if let Some(opt_15) = self.disable_output_lane.as_ref() {
662                opt_15.unparse_tokens(tokens);
663            }
664            tokens.push(PtxToken::Comma);
665            self.enable_input_d.unparse_tokens(tokens);
666            tokens.push(PtxToken::Semicolon);
667        }
668    }
669
670    impl PtxUnparser for Tcgen05MmaCtaGroupKindI81 {
671        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
672            push_opcode(tokens, "tcgen05");
673            push_directive(tokens, "mma");
674            match &self.cta_group {
675                CtaGroup::CtaGroup1 => {
676                    push_directive(tokens, "cta_group::1");
677                }
678                CtaGroup::CtaGroup2 => {
679                    push_directive(tokens, "cta_group::2");
680                }
681            }
682            push_directive(tokens, "kind::i8");
683            self.d_tmem.unparse_tokens(tokens);
684            tokens.push(PtxToken::Comma);
685            self.a_tmem.unparse_tokens(tokens);
686            tokens.push(PtxToken::Comma);
687            self.b_desc.unparse_tokens(tokens);
688            tokens.push(PtxToken::Comma);
689            self.idesc.unparse_tokens(tokens);
690            if self.disable_output_lane.is_some() {
691                tokens.push(PtxToken::Comma);
692            }
693            if let Some(opt_16) = self.disable_output_lane.as_ref() {
694                opt_16.unparse_tokens(tokens);
695            }
696            tokens.push(PtxToken::Comma);
697            self.enable_input_d.unparse_tokens(tokens);
698            tokens.push(PtxToken::Semicolon);
699        }
700    }
701}
702
703pub mod section_5 {
704    use super::*;
705    use crate::r#type::instruction::tcgen05_mma::section_5::*;
706
707    impl PtxUnparser for Tcgen05MmaCtaGroupKindI8CollectorUsage {
708        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
709            push_opcode(tokens, "tcgen05");
710            push_directive(tokens, "mma");
711            match &self.cta_group {
712                CtaGroup::CtaGroup1 => {
713                    push_directive(tokens, "cta_group::1");
714                }
715                CtaGroup::CtaGroup2 => {
716                    push_directive(tokens, "cta_group::2");
717                }
718            }
719            push_directive(tokens, "kind::i8");
720            match &self.collector_usage {
721                CollectorUsage::CollectorBufferOp(_, n1, n2) => {
722                    let mut combined = String::new();
723                    combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
724                    combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
725                    tokens.push(PtxToken::Dot);
726                    tokens.push(PtxToken::Identifier(
727                        format!("{}{}", "collector", combined).into(),
728                    ));
729                }
730            }
731            self.d_tmem.unparse_tokens(tokens);
732            tokens.push(PtxToken::Comma);
733            self.a_desc.unparse_tokens(tokens);
734            tokens.push(PtxToken::Comma);
735            self.b_desc.unparse_tokens(tokens);
736            tokens.push(PtxToken::Comma);
737            self.idesc.unparse_tokens(tokens);
738            if self.disable_output_lane.is_some() {
739                tokens.push(PtxToken::Comma);
740            }
741            if let Some(opt_17) = self.disable_output_lane.as_ref() {
742                opt_17.unparse_tokens(tokens);
743            }
744            tokens.push(PtxToken::Comma);
745            self.enable_input_d.unparse_tokens(tokens);
746            tokens.push(PtxToken::Semicolon);
747        }
748    }
749
750    impl PtxUnparser for Tcgen05MmaCtaGroupKindI8AshiftCollectorUsage {
751        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
752            push_opcode(tokens, "tcgen05");
753            push_directive(tokens, "mma");
754            match &self.cta_group {
755                CtaGroup::CtaGroup1 => {
756                    push_directive(tokens, "cta_group::1");
757                }
758                CtaGroup::CtaGroup2 => {
759                    push_directive(tokens, "cta_group::2");
760                }
761            }
762            push_directive(tokens, "kind::i8");
763            push_directive(tokens, "ashift");
764            if let Some(collector_usage_18) = self.collector_usage.as_ref() {
765                match collector_usage_18 {
766                    CollectorUsage::CollectorBufferOp(_, n1, n2) => {
767                        let mut combined = String::new();
768                        combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
769                        combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
770                        tokens.push(PtxToken::Dot);
771                        tokens.push(PtxToken::Identifier(
772                            format!("{}{}", "collector", combined).into(),
773                        ));
774                    }
775                }
776            }
777            self.d_tmem.unparse_tokens(tokens);
778            tokens.push(PtxToken::Comma);
779            self.a_tmem.unparse_tokens(tokens);
780            tokens.push(PtxToken::Comma);
781            self.b_desc.unparse_tokens(tokens);
782            tokens.push(PtxToken::Comma);
783            self.idesc.unparse_tokens(tokens);
784            if self.disable_output_lane.is_some() {
785                tokens.push(PtxToken::Comma);
786            }
787            if let Some(opt_19) = self.disable_output_lane.as_ref() {
788                opt_19.unparse_tokens(tokens);
789            }
790            tokens.push(PtxToken::Comma);
791            self.enable_input_d.unparse_tokens(tokens);
792            tokens.push(PtxToken::Semicolon);
793        }
794    }
795
796    impl PtxUnparser for Tcgen05MmaCtaGroupKindI8AshiftCollectorUsage1 {
797        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
798            push_opcode(tokens, "tcgen05");
799            push_directive(tokens, "mma");
800            match &self.cta_group {
801                CtaGroup::CtaGroup1 => {
802                    push_directive(tokens, "cta_group::1");
803                }
804                CtaGroup::CtaGroup2 => {
805                    push_directive(tokens, "cta_group::2");
806                }
807            }
808            push_directive(tokens, "kind::i8");
809            if self.ashift {
810                push_directive(tokens, "ashift");
811            }
812            match &self.collector_usage {
813                CollectorUsage::CollectorBufferOp(_, n1, n2) => {
814                    let mut combined = String::new();
815                    combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
816                    combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
817                    tokens.push(PtxToken::Dot);
818                    tokens.push(PtxToken::Identifier(
819                        format!("{}{}", "collector", combined).into(),
820                    ));
821                }
822            }
823            self.d_tmem.unparse_tokens(tokens);
824            tokens.push(PtxToken::Comma);
825            self.a_tmem.unparse_tokens(tokens);
826            tokens.push(PtxToken::Comma);
827            self.b_desc.unparse_tokens(tokens);
828            tokens.push(PtxToken::Comma);
829            self.idesc.unparse_tokens(tokens);
830            if self.disable_output_lane.is_some() {
831                tokens.push(PtxToken::Comma);
832            }
833            if let Some(opt_20) = self.disable_output_lane.as_ref() {
834                opt_20.unparse_tokens(tokens);
835            }
836            tokens.push(PtxToken::Comma);
837            self.enable_input_d.unparse_tokens(tokens);
838            tokens.push(PtxToken::Semicolon);
839        }
840    }
841}
ptx_parser/unparser/instruction/tcgen05_mma.rs

ptx_parser/unparser/instruction/
tcgen05_mma.rs
