ptx_parser/unparser/instruction/
tcgen05_mma_ws.rs

1//! Original PTX specification:
2//!
3//! // 1. Floating-point type without block scaling:
4//! tcgen05.mma.ws.cta_group::1.kind{.collector_usage}    [d-tmem],  a-desc,  b-desc,  idesc,
5//! enable-input-d {, zero-column-mask-desc };
6//! tcgen05.mma.ws.cta_group::1.kind{.collector_usage}    [d-tmem], [a-tmem], b-desc, idesc,
7//! enable-input-d {, zero-column-mask-desc };
8//! .kind = { .kind::f16, .kind::tf32, .kind::f8f6f4 };
9//! ----------------------------------------------------------------------------------
10//! // 2. Integer type:
11//! tcgen05.mma.ws.cta_group::1.kind::i8{.collector_usage} [d-tmem],  a-desc,  b-desc, idesc,
12//! enable-input-d {, zero-column-mask-desc};
13//! tcgen05.mma.ws.cta_group::1.kind::i8{.collector_usage} [d-tmem], [a-tmem], b-desc, idesc,
14//! enable-input-d {, zero-column-mask-desc};
15//! .collector_usage = { .collector::buffer::op };
16//! ::buffer = { ::b0, ::b1, ::b2, ::b3 };
17//! ::op   = { ::fill, ::use, ::lastuse, ::discard};
18
19#![allow(unused)]
20
21use crate::lexer::PtxToken;
22use crate::unparser::{PtxUnparser, common::*};
23
24pub mod section_0 {
25    use super::*;
26    use crate::r#type::instruction::tcgen05_mma_ws::section_0::*;
27
28    impl PtxUnparser for Tcgen05MmaWsCtaGroup1KindCollectorUsage {
29        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
30            push_opcode(tokens, "tcgen05");
31                    push_directive(tokens, "mma");
32                    push_directive(tokens, "ws");
33                    push_directive(tokens, "cta_group::1");
34                    match &self.kind {
35                            Kind::KindF8f6f4 => {
36                                    push_directive(tokens, "kind::f8f6f4");
37                            }
38                            Kind::KindTf32 => {
39                                    push_directive(tokens, "kind::tf32");
40                            }
41                            Kind::KindF16 => {
42                                    push_directive(tokens, "kind::f16");
43                            }
44                    }
45                    if self.collector_usage {
46                            push_directive(tokens, "collector_usage");
47                    }
48                    self.d_tmem.unparse_tokens(tokens);
49            tokens.push(PtxToken::Comma);
50                    self.a_desc.unparse_tokens(tokens);
51            tokens.push(PtxToken::Comma);
52                    self.b_desc.unparse_tokens(tokens);
53            tokens.push(PtxToken::Comma);
54                    self.idesc.unparse_tokens(tokens);
55            tokens.push(PtxToken::Comma);
56                    self.enable_input_d.unparse_tokens(tokens);
57            if self.zero_column_mask_desc.is_some() { tokens.push(PtxToken::Comma); }
58                    if let Some(opt_0) = self.zero_column_mask_desc.as_ref() {
59                        opt_0.unparse_tokens(tokens);
60                    }
61            tokens.push(PtxToken::Semicolon);
62        }
63    }
64
65    impl PtxUnparser for Tcgen05MmaWsCtaGroup1KindCollectorUsage1 {
66        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
67            push_opcode(tokens, "tcgen05");
68                    push_directive(tokens, "mma");
69                    push_directive(tokens, "ws");
70                    push_directive(tokens, "cta_group::1");
71                    match &self.kind {
72                            Kind::KindF8f6f4 => {
73                                    push_directive(tokens, "kind::f8f6f4");
74                            }
75                            Kind::KindTf32 => {
76                                    push_directive(tokens, "kind::tf32");
77                            }
78                            Kind::KindF16 => {
79                                    push_directive(tokens, "kind::f16");
80                            }
81                    }
82                    if self.collector_usage {
83                            push_directive(tokens, "collector_usage");
84                    }
85                    self.d_tmem.unparse_tokens(tokens);
86            tokens.push(PtxToken::Comma);
87                    self.a_tmem.unparse_tokens(tokens);
88            tokens.push(PtxToken::Comma);
89                    self.b_desc.unparse_tokens(tokens);
90            tokens.push(PtxToken::Comma);
91                    self.idesc.unparse_tokens(tokens);
92            tokens.push(PtxToken::Comma);
93                    self.enable_input_d.unparse_tokens(tokens);
94            if self.zero_column_mask_desc.is_some() { tokens.push(PtxToken::Comma); }
95                    if let Some(opt_1) = self.zero_column_mask_desc.as_ref() {
96                        opt_1.unparse_tokens(tokens);
97                    }
98            tokens.push(PtxToken::Semicolon);
99        }
100    }
101
102}
103
104pub mod section_1 {
105    use super::*;
106    use crate::r#type::instruction::tcgen05_mma_ws::section_1::*;
107
108    impl PtxUnparser for Tcgen05MmaWsCtaGroup1KindI8CollectorUsage {
109        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
110            push_opcode(tokens, "tcgen05");
111                    push_directive(tokens, "mma");
112                    push_directive(tokens, "ws");
113                    push_directive(tokens, "cta_group::1");
114                    push_directive(tokens, "kind::i8");
115                    if let Some(collector_usage_2) = self.collector_usage.as_ref() {
116                            match collector_usage_2 {
117                                    CollectorUsage::CollectorBufferOp(_, n1, n2) => {
118                                            let mut combined = String::new();
119                                            combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
120                                            combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
121                                            tokens.push(PtxToken::Dot);
122                                            tokens.push(PtxToken::Identifier(format!("{}{}", "collector", combined).into()));
123                                    }
124                            }
125                    }
126                    self.d_tmem.unparse_tokens(tokens);
127            tokens.push(PtxToken::Comma);
128                    self.a_desc.unparse_tokens(tokens);
129            tokens.push(PtxToken::Comma);
130                    self.b_desc.unparse_tokens(tokens);
131            tokens.push(PtxToken::Comma);
132                    self.idesc.unparse_tokens(tokens);
133            tokens.push(PtxToken::Comma);
134                    self.enable_input_d.unparse_tokens(tokens);
135            if self.zero_column_mask_desc.is_some() { tokens.push(PtxToken::Comma); }
136                    if let Some(opt_3) = self.zero_column_mask_desc.as_ref() {
137                        opt_3.unparse_tokens(tokens);
138                    }
139            tokens.push(PtxToken::Semicolon);
140        }
141    }
142
143    impl PtxUnparser for Tcgen05MmaWsCtaGroup1KindI8CollectorUsage1 {
144        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
145            push_opcode(tokens, "tcgen05");
146                    push_directive(tokens, "mma");
147                    push_directive(tokens, "ws");
148                    push_directive(tokens, "cta_group::1");
149                    push_directive(tokens, "kind::i8");
150                    if let Some(collector_usage_4) = self.collector_usage.as_ref() {
151                            match collector_usage_4 {
152                                    CollectorUsage::CollectorBufferOp(_, n1, n2) => {
153                                            let mut combined = String::new();
154                                            combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
155                                            combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
156                                            tokens.push(PtxToken::Dot);
157                                            tokens.push(PtxToken::Identifier(format!("{}{}", "collector", combined).into()));
158                                    }
159                            }
160                    }
161                    self.d_tmem.unparse_tokens(tokens);
162            tokens.push(PtxToken::Comma);
163                    self.a_tmem.unparse_tokens(tokens);
164            tokens.push(PtxToken::Comma);
165                    self.b_desc.unparse_tokens(tokens);
166            tokens.push(PtxToken::Comma);
167                    self.idesc.unparse_tokens(tokens);
168            tokens.push(PtxToken::Comma);
169                    self.enable_input_d.unparse_tokens(tokens);
170            if self.zero_column_mask_desc.is_some() { tokens.push(PtxToken::Comma); }
171                    if let Some(opt_5) = self.zero_column_mask_desc.as_ref() {
172                        opt_5.unparse_tokens(tokens);
173                    }
174            tokens.push(PtxToken::Semicolon);
175        }
176    }
177
178}
179
ptx_parser/unparser/instruction/tcgen05_mma_ws.rs

ptx_parser/unparser/instruction/
tcgen05_mma_ws.rs
