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

ptx_parser/unparser/instruction/
tcgen05_mma_ws.rs
