ptx_parser/parser/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::parser::{
22 PtxParseError, PtxParser, PtxTokenStream, Span,
23 util::{
24 between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
25 pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
26 },
27};
28use crate::r#type::common::*;
29use crate::{alt, ok, seq_n};
30
31pub mod section_0 {
32 use super::*;
33 use crate::r#type::instruction::tcgen05_mma_ws::section_0::*;
34
35 // ============================================================================
36 // Generated enum parsers
37 // ============================================================================
38
39 impl PtxParser for Kind {
40 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
41 alt!(
42 map(string_p(".kind::f8f6f4"), |_, _span| Kind::KindF8f6f4),
43 map(string_p(".kind::tf32"), |_, _span| Kind::KindTf32),
44 map(string_p(".kind::f16"), |_, _span| Kind::KindF16)
45 )
46 }
47 }
48
49 impl PtxParser for Tcgen05MmaWsCtaGroup1KindCollectorUsage {
50 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
51 try_map(
52 seq_n!(
53 string_p("tcgen05"),
54 string_p(".mma"),
55 string_p(".ws"),
56 string_p(".cta_group::1"),
57 Kind::parse(),
58 map(optional(string_p(".collector_usage")), |value, _| value
59 .is_some()),
60 AddressOperand::parse(),
61 comma_p(),
62 GeneralOperand::parse(),
63 comma_p(),
64 GeneralOperand::parse(),
65 comma_p(),
66 GeneralOperand::parse(),
67 comma_p(),
68 GeneralOperand::parse(),
69 map(
70 optional(seq_n!(comma_p(), GeneralOperand::parse())),
71 |value, _| value.map(|(_, operand)| operand)
72 ),
73 semicolon_p()
74 ),
75 |(
76 _,
77 mma,
78 ws,
79 cta_group_1,
80 kind,
81 collector_usage,
82 d_tmem,
83 _,
84 a_desc,
85 _,
86 b_desc,
87 _,
88 idesc,
89 _,
90 enable_input_d,
91 zero_column_mask_desc,
92 _,
93 ),
94 span| {
95 ok!(Tcgen05MmaWsCtaGroup1KindCollectorUsage {
96 mma = mma,
97 ws = ws,
98 cta_group_1 = cta_group_1,
99 kind = kind,
100 collector_usage = collector_usage,
101 d_tmem = d_tmem,
102 a_desc = a_desc,
103 b_desc = b_desc,
104 idesc = idesc,
105 enable_input_d = enable_input_d,
106 zero_column_mask_desc = zero_column_mask_desc,
107
108 })
109 },
110 )
111 }
112 }
113
114 impl PtxParser for Tcgen05MmaWsCtaGroup1KindCollectorUsage1 {
115 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
116 try_map(
117 seq_n!(
118 string_p("tcgen05"),
119 string_p(".mma"),
120 string_p(".ws"),
121 string_p(".cta_group::1"),
122 Kind::parse(),
123 map(optional(string_p(".collector_usage")), |value, _| value
124 .is_some()),
125 AddressOperand::parse(),
126 comma_p(),
127 AddressOperand::parse(),
128 comma_p(),
129 GeneralOperand::parse(),
130 comma_p(),
131 GeneralOperand::parse(),
132 comma_p(),
133 GeneralOperand::parse(),
134 map(
135 optional(seq_n!(comma_p(), GeneralOperand::parse())),
136 |value, _| value.map(|(_, operand)| operand)
137 ),
138 semicolon_p()
139 ),
140 |(
141 _,
142 mma,
143 ws,
144 cta_group_1,
145 kind,
146 collector_usage,
147 d_tmem,
148 _,
149 a_tmem,
150 _,
151 b_desc,
152 _,
153 idesc,
154 _,
155 enable_input_d,
156 zero_column_mask_desc,
157 _,
158 ),
159 span| {
160 ok!(Tcgen05MmaWsCtaGroup1KindCollectorUsage1 {
161 mma = mma,
162 ws = ws,
163 cta_group_1 = cta_group_1,
164 kind = kind,
165 collector_usage = collector_usage,
166 d_tmem = d_tmem,
167 a_tmem = a_tmem,
168 b_desc = b_desc,
169 idesc = idesc,
170 enable_input_d = enable_input_d,
171 zero_column_mask_desc = zero_column_mask_desc,
172
173 })
174 },
175 )
176 }
177 }
178}
179
180pub mod section_1 {
181 use super::*;
182 use crate::r#type::instruction::tcgen05_mma_ws::section_1::*;
183
184 // ============================================================================
185 // Generated enum parsers
186 // ============================================================================
187
188 impl PtxParser for CollectorUsage {
189 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
190 alt!(map(
191 |stream| {
192 stream.try_with_span(|stream| {
193 stream.with_partial_token_mode(|stream| {
194 stream.expect_string(".collector")?;
195 let part0 =
196 match stream.expect_strings(&["::b0", "::b1", "::b2", "::b3"])? {
197 0 => Buffer::B0,
198 1 => Buffer::B1,
199 2 => Buffer::B2,
200 3 => Buffer::B3,
201 _ => unreachable!(),
202 };
203 let part1 = match stream.expect_strings(&[
204 "::lastuse",
205 "::discard",
206 "::fill",
207 "::use",
208 ])? {
209 0 => Op::Lastuse,
210 1 => Op::Discard,
211 2 => Op::Fill,
212 3 => Op::Use,
213 _ => unreachable!(),
214 };
215 Ok(((), part0, part1))
216 })
217 })
218 },
219 |(collector, buffer, op), _span| CollectorUsage::CollectorBufferOp(
220 collector, buffer, op
221 )
222 ))
223 }
224 }
225
226 impl PtxParser for Tcgen05MmaWsCtaGroup1KindI8CollectorUsage {
227 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
228 try_map(
229 seq_n!(
230 string_p("tcgen05"),
231 string_p(".mma"),
232 string_p(".ws"),
233 string_p(".cta_group::1"),
234 string_p(".kind::i8"),
235 optional(CollectorUsage::parse()),
236 AddressOperand::parse(),
237 comma_p(),
238 GeneralOperand::parse(),
239 comma_p(),
240 GeneralOperand::parse(),
241 comma_p(),
242 GeneralOperand::parse(),
243 comma_p(),
244 GeneralOperand::parse(),
245 map(
246 optional(seq_n!(comma_p(), GeneralOperand::parse())),
247 |value, _| value.map(|(_, operand)| operand)
248 ),
249 semicolon_p()
250 ),
251 |(
252 _,
253 mma,
254 ws,
255 cta_group_1,
256 kind_i8,
257 collector_usage,
258 d_tmem,
259 _,
260 a_desc,
261 _,
262 b_desc,
263 _,
264 idesc,
265 _,
266 enable_input_d,
267 zero_column_mask_desc,
268 _,
269 ),
270 span| {
271 ok!(Tcgen05MmaWsCtaGroup1KindI8CollectorUsage {
272 mma = mma,
273 ws = ws,
274 cta_group_1 = cta_group_1,
275 kind_i8 = kind_i8,
276 collector_usage = collector_usage,
277 d_tmem = d_tmem,
278 a_desc = a_desc,
279 b_desc = b_desc,
280 idesc = idesc,
281 enable_input_d = enable_input_d,
282 zero_column_mask_desc = zero_column_mask_desc,
283
284 })
285 },
286 )
287 }
288 }
289
290 impl PtxParser for Tcgen05MmaWsCtaGroup1KindI8CollectorUsage1 {
291 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
292 try_map(
293 seq_n!(
294 string_p("tcgen05"),
295 string_p(".mma"),
296 string_p(".ws"),
297 string_p(".cta_group::1"),
298 string_p(".kind::i8"),
299 optional(CollectorUsage::parse()),
300 AddressOperand::parse(),
301 comma_p(),
302 AddressOperand::parse(),
303 comma_p(),
304 GeneralOperand::parse(),
305 comma_p(),
306 GeneralOperand::parse(),
307 comma_p(),
308 GeneralOperand::parse(),
309 map(
310 optional(seq_n!(comma_p(), GeneralOperand::parse())),
311 |value, _| value.map(|(_, operand)| operand)
312 ),
313 semicolon_p()
314 ),
315 |(
316 _,
317 mma,
318 ws,
319 cta_group_1,
320 kind_i8,
321 collector_usage,
322 d_tmem,
323 _,
324 a_tmem,
325 _,
326 b_desc,
327 _,
328 idesc,
329 _,
330 enable_input_d,
331 zero_column_mask_desc,
332 _,
333 ),
334 span| {
335 ok!(Tcgen05MmaWsCtaGroup1KindI8CollectorUsage1 {
336 mma = mma,
337 ws = ws,
338 cta_group_1 = cta_group_1,
339 kind_i8 = kind_i8,
340 collector_usage = collector_usage,
341 d_tmem = d_tmem,
342 a_tmem = a_tmem,
343 b_desc = b_desc,
344 idesc = idesc,
345 enable_input_d = enable_input_d,
346 zero_column_mask_desc = zero_column_mask_desc,
347
348 })
349 },
350 )
351 }
352 }
353}