ptx_parser/unparser/instruction/
tcgen05_mma.rs

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

ptx_parser/unparser/instruction/
tcgen05_mma.rs
