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            self.unparse_tokens_mode(tokens, false);
31        }
32        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
33            push_opcode(tokens, "tcgen05");
34            push_directive(tokens, "mma");
35            push_directive(tokens, "ws");
36            push_directive(tokens, "cta_group::1");
37            match &self.kind {
38                Kind::KindF8f6f4 => {
39                    push_directive(tokens, "kind::f8f6f4");
40                }
41                Kind::KindTf32 => {
42                    push_directive(tokens, "kind::tf32");
43                }
44                Kind::KindF16 => {
45                    push_directive(tokens, "kind::f16");
46                }
47            }
48            if self.collector_usage {
49                push_directive(tokens, "collector_usage");
50            }
51            if spaced {
52                tokens.push(PtxToken::Space);
53            }
54            self.d_tmem.unparse_tokens_mode(tokens, spaced);
55            tokens.push(PtxToken::Comma);
56            if spaced {
57                tokens.push(PtxToken::Space);
58            }
59            self.a_desc.unparse_tokens_mode(tokens, spaced);
60            tokens.push(PtxToken::Comma);
61            if spaced {
62                tokens.push(PtxToken::Space);
63            }
64            self.b_desc.unparse_tokens_mode(tokens, spaced);
65            tokens.push(PtxToken::Comma);
66            if spaced {
67                tokens.push(PtxToken::Space);
68            }
69            self.idesc.unparse_tokens_mode(tokens, spaced);
70            tokens.push(PtxToken::Comma);
71            if spaced {
72                tokens.push(PtxToken::Space);
73            }
74            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
75            if self.zero_column_mask_desc.is_some() {
76                tokens.push(PtxToken::Comma);
77            }
78            if let Some(opt_0) = self.zero_column_mask_desc.as_ref() {
79                if spaced {
80                    tokens.push(PtxToken::Space);
81                }
82                opt_0.unparse_tokens_mode(tokens, spaced);
83            }
84            tokens.push(PtxToken::Semicolon);
85            if spaced {
86                tokens.push(PtxToken::Newline);
87            }
88        }
89    }
90
91    impl PtxUnparser for Tcgen05MmaWsCtaGroup1KindCollectorUsage1 {
92        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
93            self.unparse_tokens_mode(tokens, false);
94        }
95        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
96            push_opcode(tokens, "tcgen05");
97            push_directive(tokens, "mma");
98            push_directive(tokens, "ws");
99            push_directive(tokens, "cta_group::1");
100            match &self.kind {
101                Kind::KindF8f6f4 => {
102                    push_directive(tokens, "kind::f8f6f4");
103                }
104                Kind::KindTf32 => {
105                    push_directive(tokens, "kind::tf32");
106                }
107                Kind::KindF16 => {
108                    push_directive(tokens, "kind::f16");
109                }
110            }
111            if self.collector_usage {
112                push_directive(tokens, "collector_usage");
113            }
114            if spaced {
115                tokens.push(PtxToken::Space);
116            }
117            self.d_tmem.unparse_tokens_mode(tokens, spaced);
118            tokens.push(PtxToken::Comma);
119            if spaced {
120                tokens.push(PtxToken::Space);
121            }
122            self.a_tmem.unparse_tokens_mode(tokens, spaced);
123            tokens.push(PtxToken::Comma);
124            if spaced {
125                tokens.push(PtxToken::Space);
126            }
127            self.b_desc.unparse_tokens_mode(tokens, spaced);
128            tokens.push(PtxToken::Comma);
129            if spaced {
130                tokens.push(PtxToken::Space);
131            }
132            self.idesc.unparse_tokens_mode(tokens, spaced);
133            tokens.push(PtxToken::Comma);
134            if spaced {
135                tokens.push(PtxToken::Space);
136            }
137            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
138            if self.zero_column_mask_desc.is_some() {
139                tokens.push(PtxToken::Comma);
140            }
141            if let Some(opt_1) = self.zero_column_mask_desc.as_ref() {
142                if spaced {
143                    tokens.push(PtxToken::Space);
144                }
145                opt_1.unparse_tokens_mode(tokens, spaced);
146            }
147            tokens.push(PtxToken::Semicolon);
148            if spaced {
149                tokens.push(PtxToken::Newline);
150            }
151        }
152    }
153}
154
155pub mod section_1 {
156    use super::*;
157    use crate::r#type::instruction::tcgen05_mma_ws::section_1::*;
158
159    impl PtxUnparser for Tcgen05MmaWsCtaGroup1KindI8CollectorUsage {
160        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
161            self.unparse_tokens_mode(tokens, false);
162        }
163        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
164            push_opcode(tokens, "tcgen05");
165            push_directive(tokens, "mma");
166            push_directive(tokens, "ws");
167            push_directive(tokens, "cta_group::1");
168            push_directive(tokens, "kind::i8");
169            if let Some(collector_usage_2) = self.collector_usage.as_ref() {
170                match collector_usage_2 {
171                    CollectorUsage::CollectorBufferOp(_, n1, n2) => {
172                        let mut combined = String::new();
173                        combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
174                        combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
175                        tokens.push(PtxToken::Dot);
176                        tokens.push(PtxToken::Identifier(
177                            format!("{}{}", "collector", combined).into(),
178                        ));
179                    }
180                }
181            }
182            if spaced {
183                tokens.push(PtxToken::Space);
184            }
185            self.d_tmem.unparse_tokens_mode(tokens, spaced);
186            tokens.push(PtxToken::Comma);
187            if spaced {
188                tokens.push(PtxToken::Space);
189            }
190            self.a_desc.unparse_tokens_mode(tokens, spaced);
191            tokens.push(PtxToken::Comma);
192            if spaced {
193                tokens.push(PtxToken::Space);
194            }
195            self.b_desc.unparse_tokens_mode(tokens, spaced);
196            tokens.push(PtxToken::Comma);
197            if spaced {
198                tokens.push(PtxToken::Space);
199            }
200            self.idesc.unparse_tokens_mode(tokens, spaced);
201            tokens.push(PtxToken::Comma);
202            if spaced {
203                tokens.push(PtxToken::Space);
204            }
205            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
206            if self.zero_column_mask_desc.is_some() {
207                tokens.push(PtxToken::Comma);
208            }
209            if let Some(opt_3) = self.zero_column_mask_desc.as_ref() {
210                if spaced {
211                    tokens.push(PtxToken::Space);
212                }
213                opt_3.unparse_tokens_mode(tokens, spaced);
214            }
215            tokens.push(PtxToken::Semicolon);
216            if spaced {
217                tokens.push(PtxToken::Newline);
218            }
219        }
220    }
221
222    impl PtxUnparser for Tcgen05MmaWsCtaGroup1KindI8CollectorUsage1 {
223        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
224            self.unparse_tokens_mode(tokens, false);
225        }
226        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
227            push_opcode(tokens, "tcgen05");
228            push_directive(tokens, "mma");
229            push_directive(tokens, "ws");
230            push_directive(tokens, "cta_group::1");
231            push_directive(tokens, "kind::i8");
232            if let Some(collector_usage_4) = self.collector_usage.as_ref() {
233                match collector_usage_4 {
234                    CollectorUsage::CollectorBufferOp(_, n1, n2) => {
235                        let mut combined = String::new();
236                        combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
237                        combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
238                        tokens.push(PtxToken::Dot);
239                        tokens.push(PtxToken::Identifier(
240                            format!("{}{}", "collector", combined).into(),
241                        ));
242                    }
243                }
244            }
245            if spaced {
246                tokens.push(PtxToken::Space);
247            }
248            self.d_tmem.unparse_tokens_mode(tokens, spaced);
249            tokens.push(PtxToken::Comma);
250            if spaced {
251                tokens.push(PtxToken::Space);
252            }
253            self.a_tmem.unparse_tokens_mode(tokens, spaced);
254            tokens.push(PtxToken::Comma);
255            if spaced {
256                tokens.push(PtxToken::Space);
257            }
258            self.b_desc.unparse_tokens_mode(tokens, spaced);
259            tokens.push(PtxToken::Comma);
260            if spaced {
261                tokens.push(PtxToken::Space);
262            }
263            self.idesc.unparse_tokens_mode(tokens, spaced);
264            tokens.push(PtxToken::Comma);
265            if spaced {
266                tokens.push(PtxToken::Space);
267            }
268            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
269            if self.zero_column_mask_desc.is_some() {
270                tokens.push(PtxToken::Comma);
271            }
272            if let Some(opt_5) = self.zero_column_mask_desc.as_ref() {
273                if spaced {
274                    tokens.push(PtxToken::Space);
275                }
276                opt_5.unparse_tokens_mode(tokens, spaced);
277            }
278            tokens.push(PtxToken::Semicolon);
279            if spaced {
280                tokens.push(PtxToken::Newline);
281            }
282        }
283    }
284}
ptx_parser/unparser/instruction/tcgen05_mma_ws.rs

ptx_parser/unparser/instruction/
tcgen05_mma_ws.rs
