ptx_parser/unparser/instruction/
tcgen05_mma_ws_sp.rs

1//! Original PTX specification:
2//!
3//! // 1. Floating-point type without block scaling:
4//! tcgen05.mma.ws.sp.cta_group::1.kind{.collector_usage} [d-tmem],  a-desc,  b-desc,
5//! [sp-meta-tmem] ,  idesc,
6//! enable-input-d {, zero-column-mask-desc};
7//! tcgen05.mma.ws.sp.cta_group::1.kind{.collector_usage} [d-tmem], [a-tmem], b-desc,
8//! [sp-meta-tmem] , idesc,
9//! enable-input-d {, zero-column-mask-desc};
10//! .kind = { .kind::f16, .kind::tf32, .kind::f8f6f4 };
11//! ------------------------------------------------------------------
12//! // 2. Integer type:
13//! tcgen05.mma.ws.sp.cta_group::1.kind::i8{.collector_usage} [d-tmem], a-desc, b-desc,
14//! [sp-meta-tmem] , idesc,
15//! enable-input-d {, zero-column-mask-desc};
16//! tcgen05.mma.ws.sp.cta_group::1.kind::i8{.collector_usage} [d-tmem], [a-tmem], b-desc,
17//! [sp-meta-tmem] , idesc,
18//! enable-input-d {, zero-column-mask-desc};
19//! .collector_usage = { .collector::buffer::op };
20//! ::buffer = { ::b0, ::b1, ::b2, ::b3 };
21//! ::op   = { ::fill, ::use, ::lastuse, ::discard};
22
23#![allow(unused)]
24
25use crate::lexer::PtxToken;
26use crate::unparser::{PtxUnparser, common::*};
27
28pub mod section_0 {
29    use super::*;
30    use crate::r#type::instruction::tcgen05_mma_ws_sp::section_0::*;
31
32    impl PtxUnparser for Tcgen05MmaWsSpCtaGroup1KindCollectorUsage {
33        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
34            self.unparse_tokens_mode(tokens, false);
35        }
36        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
37            push_opcode(tokens, "tcgen05");
38            push_directive(tokens, "mma");
39            push_directive(tokens, "ws");
40            push_directive(tokens, "sp");
41            push_directive(tokens, "cta_group::1");
42            match &self.kind {
43                Kind::KindF8f6f4 => {
44                    push_directive(tokens, "kind::f8f6f4");
45                }
46                Kind::KindTf32 => {
47                    push_directive(tokens, "kind::tf32");
48                }
49                Kind::KindF16 => {
50                    push_directive(tokens, "kind::f16");
51                }
52            }
53            if self.collector_usage {
54                push_directive(tokens, "collector_usage");
55            }
56            if spaced {
57                tokens.push(PtxToken::Space);
58            }
59            self.d_tmem.unparse_tokens_mode(tokens, spaced);
60            tokens.push(PtxToken::Comma);
61            if spaced {
62                tokens.push(PtxToken::Space);
63            }
64            self.a_desc.unparse_tokens_mode(tokens, spaced);
65            tokens.push(PtxToken::Comma);
66            if spaced {
67                tokens.push(PtxToken::Space);
68            }
69            self.b_desc.unparse_tokens_mode(tokens, spaced);
70            tokens.push(PtxToken::Comma);
71            if spaced {
72                tokens.push(PtxToken::Space);
73            }
74            self.sp_meta_tmem.unparse_tokens_mode(tokens, spaced);
75            tokens.push(PtxToken::Comma);
76            if spaced {
77                tokens.push(PtxToken::Space);
78            }
79            self.idesc.unparse_tokens_mode(tokens, spaced);
80            tokens.push(PtxToken::Comma);
81            if spaced {
82                tokens.push(PtxToken::Space);
83            }
84            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
85            if self.zero_column_mask_desc.is_some() {
86                tokens.push(PtxToken::Comma);
87            }
88            if let Some(opt_0) = self.zero_column_mask_desc.as_ref() {
89                if spaced {
90                    tokens.push(PtxToken::Space);
91                }
92                opt_0.unparse_tokens_mode(tokens, spaced);
93            }
94            tokens.push(PtxToken::Semicolon);
95            if spaced {
96                tokens.push(PtxToken::Newline);
97            }
98        }
99    }
100
101    impl PtxUnparser for Tcgen05MmaWsSpCtaGroup1KindCollectorUsage1 {
102        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
103            self.unparse_tokens_mode(tokens, false);
104        }
105        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
106            push_opcode(tokens, "tcgen05");
107            push_directive(tokens, "mma");
108            push_directive(tokens, "ws");
109            push_directive(tokens, "sp");
110            push_directive(tokens, "cta_group::1");
111            match &self.kind {
112                Kind::KindF8f6f4 => {
113                    push_directive(tokens, "kind::f8f6f4");
114                }
115                Kind::KindTf32 => {
116                    push_directive(tokens, "kind::tf32");
117                }
118                Kind::KindF16 => {
119                    push_directive(tokens, "kind::f16");
120                }
121            }
122            if self.collector_usage {
123                push_directive(tokens, "collector_usage");
124            }
125            if spaced {
126                tokens.push(PtxToken::Space);
127            }
128            self.d_tmem.unparse_tokens_mode(tokens, spaced);
129            tokens.push(PtxToken::Comma);
130            if spaced {
131                tokens.push(PtxToken::Space);
132            }
133            self.a_tmem.unparse_tokens_mode(tokens, spaced);
134            tokens.push(PtxToken::Comma);
135            if spaced {
136                tokens.push(PtxToken::Space);
137            }
138            self.b_desc.unparse_tokens_mode(tokens, spaced);
139            tokens.push(PtxToken::Comma);
140            if spaced {
141                tokens.push(PtxToken::Space);
142            }
143            self.sp_meta_tmem.unparse_tokens_mode(tokens, spaced);
144            tokens.push(PtxToken::Comma);
145            if spaced {
146                tokens.push(PtxToken::Space);
147            }
148            self.idesc.unparse_tokens_mode(tokens, spaced);
149            tokens.push(PtxToken::Comma);
150            if spaced {
151                tokens.push(PtxToken::Space);
152            }
153            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
154            if self.zero_column_mask_desc.is_some() {
155                tokens.push(PtxToken::Comma);
156            }
157            if let Some(opt_1) = self.zero_column_mask_desc.as_ref() {
158                if spaced {
159                    tokens.push(PtxToken::Space);
160                }
161                opt_1.unparse_tokens_mode(tokens, spaced);
162            }
163            tokens.push(PtxToken::Semicolon);
164            if spaced {
165                tokens.push(PtxToken::Newline);
166            }
167        }
168    }
169}
170
171pub mod section_1 {
172    use super::*;
173    use crate::r#type::instruction::tcgen05_mma_ws_sp::section_1::*;
174
175    impl PtxUnparser for Tcgen05MmaWsSpCtaGroup1KindI8CollectorUsage {
176        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
177            self.unparse_tokens_mode(tokens, false);
178        }
179        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
180            push_opcode(tokens, "tcgen05");
181            push_directive(tokens, "mma");
182            push_directive(tokens, "ws");
183            push_directive(tokens, "sp");
184            push_directive(tokens, "cta_group::1");
185            push_directive(tokens, "kind::i8");
186            if let Some(collector_usage_2) = self.collector_usage.as_ref() {
187                match collector_usage_2 {
188                    CollectorUsage::CollectorBufferOp(_, n1, n2) => {
189                        let mut combined = String::new();
190                        combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
191                        combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
192                        tokens.push(PtxToken::Dot);
193                        tokens.push(PtxToken::Identifier(
194                            format!("{}{}", "collector", combined).into(),
195                        ));
196                    }
197                }
198            }
199            if spaced {
200                tokens.push(PtxToken::Space);
201            }
202            self.d_tmem.unparse_tokens_mode(tokens, spaced);
203            tokens.push(PtxToken::Comma);
204            if spaced {
205                tokens.push(PtxToken::Space);
206            }
207            self.a_desc.unparse_tokens_mode(tokens, spaced);
208            tokens.push(PtxToken::Comma);
209            if spaced {
210                tokens.push(PtxToken::Space);
211            }
212            self.b_desc.unparse_tokens_mode(tokens, spaced);
213            tokens.push(PtxToken::Comma);
214            if spaced {
215                tokens.push(PtxToken::Space);
216            }
217            self.sp_meta_tmem.unparse_tokens_mode(tokens, spaced);
218            tokens.push(PtxToken::Comma);
219            if spaced {
220                tokens.push(PtxToken::Space);
221            }
222            self.idesc.unparse_tokens_mode(tokens, spaced);
223            tokens.push(PtxToken::Comma);
224            if spaced {
225                tokens.push(PtxToken::Space);
226            }
227            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
228            if self.zero_column_mask_desc.is_some() {
229                tokens.push(PtxToken::Comma);
230            }
231            if let Some(opt_3) = self.zero_column_mask_desc.as_ref() {
232                if spaced {
233                    tokens.push(PtxToken::Space);
234                }
235                opt_3.unparse_tokens_mode(tokens, spaced);
236            }
237            tokens.push(PtxToken::Semicolon);
238            if spaced {
239                tokens.push(PtxToken::Newline);
240            }
241        }
242    }
243
244    impl PtxUnparser for Tcgen05MmaWsSpCtaGroup1KindI8CollectorUsage1 {
245        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
246            self.unparse_tokens_mode(tokens, false);
247        }
248        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
249            push_opcode(tokens, "tcgen05");
250            push_directive(tokens, "mma");
251            push_directive(tokens, "ws");
252            push_directive(tokens, "sp");
253            push_directive(tokens, "cta_group::1");
254            push_directive(tokens, "kind::i8");
255            if let Some(collector_usage_4) = self.collector_usage.as_ref() {
256                match collector_usage_4 {
257                    CollectorUsage::CollectorBufferOp(_, n1, n2) => {
258                        let mut combined = String::new();
259                        combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
260                        combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
261                        tokens.push(PtxToken::Dot);
262                        tokens.push(PtxToken::Identifier(
263                            format!("{}{}", "collector", combined).into(),
264                        ));
265                    }
266                }
267            }
268            if spaced {
269                tokens.push(PtxToken::Space);
270            }
271            self.d_tmem.unparse_tokens_mode(tokens, spaced);
272            tokens.push(PtxToken::Comma);
273            if spaced {
274                tokens.push(PtxToken::Space);
275            }
276            self.a_tmem.unparse_tokens_mode(tokens, spaced);
277            tokens.push(PtxToken::Comma);
278            if spaced {
279                tokens.push(PtxToken::Space);
280            }
281            self.b_desc.unparse_tokens_mode(tokens, spaced);
282            tokens.push(PtxToken::Comma);
283            if spaced {
284                tokens.push(PtxToken::Space);
285            }
286            self.sp_meta_tmem.unparse_tokens_mode(tokens, spaced);
287            tokens.push(PtxToken::Comma);
288            if spaced {
289                tokens.push(PtxToken::Space);
290            }
291            self.idesc.unparse_tokens_mode(tokens, spaced);
292            tokens.push(PtxToken::Comma);
293            if spaced {
294                tokens.push(PtxToken::Space);
295            }
296            self.enable_input_d.unparse_tokens_mode(tokens, spaced);
297            if self.zero_column_mask_desc.is_some() {
298                tokens.push(PtxToken::Comma);
299            }
300            if let Some(opt_5) = self.zero_column_mask_desc.as_ref() {
301                if spaced {
302                    tokens.push(PtxToken::Space);
303                }
304                opt_5.unparse_tokens_mode(tokens, spaced);
305            }
306            tokens.push(PtxToken::Semicolon);
307            if spaced {
308                tokens.push(PtxToken::Newline);
309            }
310        }
311    }
312}
ptx_parser/unparser/instruction/tcgen05_mma_ws_sp.rs

ptx_parser/unparser/instruction/
tcgen05_mma_ws_sp.rs
