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

ptx_parser/unparser/instruction/
tcgen05_mma_sp.rs
