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