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            push_opcode(tokens, "tcgen05");
35            push_directive(tokens, "mma");
36            push_directive(tokens, "ws");
37            push_directive(tokens, "sp");
38            push_directive(tokens, "cta_group::1");
39            match &self.kind {
40                Kind::KindF8f6f4 => {
41                    push_directive(tokens, "kind::f8f6f4");
42                }
43                Kind::KindTf32 => {
44                    push_directive(tokens, "kind::tf32");
45                }
46                Kind::KindF16 => {
47                    push_directive(tokens, "kind::f16");
48                }
49            }
50            if self.collector_usage {
51                push_directive(tokens, "collector_usage");
52            }
53            self.d_tmem.unparse_tokens(tokens);
54            tokens.push(PtxToken::Comma);
55            self.a_desc.unparse_tokens(tokens);
56            tokens.push(PtxToken::Comma);
57            self.b_desc.unparse_tokens(tokens);
58            tokens.push(PtxToken::Comma);
59            self.sp_meta_tmem.unparse_tokens(tokens);
60            tokens.push(PtxToken::Comma);
61            self.idesc.unparse_tokens(tokens);
62            tokens.push(PtxToken::Comma);
63            self.enable_input_d.unparse_tokens(tokens);
64            if self.zero_column_mask_desc.is_some() {
65                tokens.push(PtxToken::Comma);
66            }
67            if let Some(opt_0) = self.zero_column_mask_desc.as_ref() {
68                opt_0.unparse_tokens(tokens);
69            }
70            tokens.push(PtxToken::Semicolon);
71        }
72    }
73
74    impl PtxUnparser for Tcgen05MmaWsSpCtaGroup1KindCollectorUsage1 {
75        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
76            push_opcode(tokens, "tcgen05");
77            push_directive(tokens, "mma");
78            push_directive(tokens, "ws");
79            push_directive(tokens, "sp");
80            push_directive(tokens, "cta_group::1");
81            match &self.kind {
82                Kind::KindF8f6f4 => {
83                    push_directive(tokens, "kind::f8f6f4");
84                }
85                Kind::KindTf32 => {
86                    push_directive(tokens, "kind::tf32");
87                }
88                Kind::KindF16 => {
89                    push_directive(tokens, "kind::f16");
90                }
91            }
92            if self.collector_usage {
93                push_directive(tokens, "collector_usage");
94            }
95            self.d_tmem.unparse_tokens(tokens);
96            tokens.push(PtxToken::Comma);
97            self.a_tmem.unparse_tokens(tokens);
98            tokens.push(PtxToken::Comma);
99            self.b_desc.unparse_tokens(tokens);
100            tokens.push(PtxToken::Comma);
101            self.sp_meta_tmem.unparse_tokens(tokens);
102            tokens.push(PtxToken::Comma);
103            self.idesc.unparse_tokens(tokens);
104            tokens.push(PtxToken::Comma);
105            self.enable_input_d.unparse_tokens(tokens);
106            if self.zero_column_mask_desc.is_some() {
107                tokens.push(PtxToken::Comma);
108            }
109            if let Some(opt_1) = self.zero_column_mask_desc.as_ref() {
110                opt_1.unparse_tokens(tokens);
111            }
112            tokens.push(PtxToken::Semicolon);
113        }
114    }
115}
116
117pub mod section_1 {
118    use super::*;
119    use crate::r#type::instruction::tcgen05_mma_ws_sp::section_1::*;
120
121    impl PtxUnparser for Tcgen05MmaWsSpCtaGroup1KindI8CollectorUsage {
122        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
123            push_opcode(tokens, "tcgen05");
124            push_directive(tokens, "mma");
125            push_directive(tokens, "ws");
126            push_directive(tokens, "sp");
127            push_directive(tokens, "cta_group::1");
128            push_directive(tokens, "kind::i8");
129            if let Some(collector_usage_2) = self.collector_usage.as_ref() {
130                match collector_usage_2 {
131                    CollectorUsage::CollectorBufferOp(_, n1, n2) => {
132                        let mut combined = String::new();
133                        combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
134                        combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
135                        tokens.push(PtxToken::Dot);
136                        tokens.push(PtxToken::Identifier(
137                            format!("{}{}", "collector", combined).into(),
138                        ));
139                    }
140                }
141            }
142            self.d_tmem.unparse_tokens(tokens);
143            tokens.push(PtxToken::Comma);
144            self.a_desc.unparse_tokens(tokens);
145            tokens.push(PtxToken::Comma);
146            self.b_desc.unparse_tokens(tokens);
147            tokens.push(PtxToken::Comma);
148            self.sp_meta_tmem.unparse_tokens(tokens);
149            tokens.push(PtxToken::Comma);
150            self.idesc.unparse_tokens(tokens);
151            tokens.push(PtxToken::Comma);
152            self.enable_input_d.unparse_tokens(tokens);
153            if self.zero_column_mask_desc.is_some() {
154                tokens.push(PtxToken::Comma);
155            }
156            if let Some(opt_3) = self.zero_column_mask_desc.as_ref() {
157                opt_3.unparse_tokens(tokens);
158            }
159            tokens.push(PtxToken::Semicolon);
160        }
161    }
162
163    impl PtxUnparser for Tcgen05MmaWsSpCtaGroup1KindI8CollectorUsage1 {
164        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
165            push_opcode(tokens, "tcgen05");
166            push_directive(tokens, "mma");
167            push_directive(tokens, "ws");
168            push_directive(tokens, "sp");
169            push_directive(tokens, "cta_group::1");
170            push_directive(tokens, "kind::i8");
171            if let Some(collector_usage_4) = self.collector_usage.as_ref() {
172                match collector_usage_4 {
173                    CollectorUsage::CollectorBufferOp(_, n1, n2) => {
174                        let mut combined = String::new();
175                        combined.push_str(format!("{:?}", n1).trim_start_matches('_'));
176                        combined.push_str(format!("{:?}", n2).trim_start_matches('_'));
177                        tokens.push(PtxToken::Dot);
178                        tokens.push(PtxToken::Identifier(
179                            format!("{}{}", "collector", combined).into(),
180                        ));
181                    }
182                }
183            }
184            self.d_tmem.unparse_tokens(tokens);
185            tokens.push(PtxToken::Comma);
186            self.a_tmem.unparse_tokens(tokens);
187            tokens.push(PtxToken::Comma);
188            self.b_desc.unparse_tokens(tokens);
189            tokens.push(PtxToken::Comma);
190            self.sp_meta_tmem.unparse_tokens(tokens);
191            tokens.push(PtxToken::Comma);
192            self.idesc.unparse_tokens(tokens);
193            tokens.push(PtxToken::Comma);
194            self.enable_input_d.unparse_tokens(tokens);
195            if self.zero_column_mask_desc.is_some() {
196                tokens.push(PtxToken::Comma);
197            }
198            if let Some(opt_5) = self.zero_column_mask_desc.as_ref() {
199                opt_5.unparse_tokens(tokens);
200            }
201            tokens.push(PtxToken::Semicolon);
202        }
203    }
204}
ptx_parser/unparser/instruction/tcgen05_mma_ws_sp.rs

ptx_parser/unparser/instruction/
tcgen05_mma_ws_sp.rs
