ptx_parser/unparser/instruction/
tcgen05_mma_sp.rs

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

ptx_parser/unparser/instruction/
tcgen05_mma_sp.rs
