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

ptx_parser/unparser/instruction/
tcgen05_mma_sp.rs
