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

ptx_parser/unparser/instruction/
tcgen05_mma_ws_sp.rs
