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::lexer::PtxToken;
22use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
23use crate::r#type::common::*;
24
25pub mod section_0 {
26    use super::*;
27    use crate::r#type::instruction::tcgen05_mma_ws::section_0::*;
28
29    // ============================================================================
30    // Generated enum parsers
31    // ============================================================================
32
33    impl PtxParser for Kind {
34        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
35            // Try KindF8f6f4
36            {
37                let saved_pos = stream.position();
38                if stream.expect_string(".kind::f8f6f4").is_ok() {
39                    return Ok(Kind::KindF8f6f4);
40                }
41                stream.set_position(saved_pos);
42            }
43            let saved_pos = stream.position();
44            // Try KindTf32
45            {
46                let saved_pos = stream.position();
47                if stream.expect_string(".kind::tf32").is_ok() {
48                    return Ok(Kind::KindTf32);
49                }
50                stream.set_position(saved_pos);
51            }
52            stream.set_position(saved_pos);
53            let saved_pos = stream.position();
54            // Try KindF16
55            {
56                let saved_pos = stream.position();
57                if stream.expect_string(".kind::f16").is_ok() {
58                    return Ok(Kind::KindF16);
59                }
60                stream.set_position(saved_pos);
61            }
62            stream.set_position(saved_pos);
63            let span = stream
64                .peek()
65                .map(|(_, s)| s.clone())
66                .unwrap_or(Span { start: 0, end: 0 });
67            let expected = &[".kind::f8f6f4", ".kind::tf32", ".kind::f16"];
68            let found = stream
69                .peek()
70                .map(|(t, _)| format!("{:?}", t))
71                .unwrap_or_else(|_| "<end of input>".to_string());
72            Err(crate::parser::unexpected_value(span, expected, found))
73        }
74    }
75
76    impl PtxParser for Tcgen05MmaWsCtaGroup1KindCollectorUsage {
77        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
78            stream.expect_string("tcgen05")?;
79            stream.expect_string(".mma")?;
80            let mma = ();
81            stream.expect_complete()?;
82            stream.expect_string(".ws")?;
83            let ws = ();
84            stream.expect_complete()?;
85            stream.expect_string(".cta_group::1")?;
86            let cta_group_1 = ();
87            stream.expect_complete()?;
88            let kind = Kind::parse(stream)?;
89            stream.expect_complete()?;
90            let saved_pos = stream.position();
91            let collector_usage = stream.expect_string(".collector_usage").is_ok();
92            if !collector_usage {
93                stream.set_position(saved_pos);
94            }
95            stream.expect_complete()?;
96            let d_tmem = AddressOperand::parse(stream)?;
97            stream.expect_complete()?;
98            stream.expect(&PtxToken::Comma)?;
99            let a_desc = GeneralOperand::parse(stream)?;
100            stream.expect_complete()?;
101            stream.expect(&PtxToken::Comma)?;
102            let b_desc = GeneralOperand::parse(stream)?;
103            stream.expect_complete()?;
104            stream.expect(&PtxToken::Comma)?;
105            let idesc = GeneralOperand::parse(stream)?;
106            stream.expect_complete()?;
107            stream.expect(&PtxToken::Comma)?;
108            let enable_input_d = GeneralOperand::parse(stream)?;
109            stream.expect_complete()?;
110            let saved_pos = stream.position();
111            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
112            if !has_comma {
113                stream.set_position(saved_pos);
114            }
115            let saved_pos = stream.position();
116            let zero_column_mask_desc = match GeneralOperand::parse(stream) {
117                Ok(val) => Some(val),
118                Err(_) => {
119                    stream.set_position(saved_pos);
120                    None
121                }
122            };
123            stream.expect_complete()?;
124            stream.expect_complete()?;
125            stream.expect(&PtxToken::Semicolon)?;
126            Ok(Tcgen05MmaWsCtaGroup1KindCollectorUsage {
127                mma,
128                ws,
129                cta_group_1,
130                kind,
131                collector_usage,
132                d_tmem,
133                a_desc,
134                b_desc,
135                idesc,
136                enable_input_d,
137                zero_column_mask_desc,
138            })
139        }
140    }
141
142    impl PtxParser for Tcgen05MmaWsCtaGroup1KindCollectorUsage1 {
143        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
144            stream.expect_string("tcgen05")?;
145            stream.expect_string(".mma")?;
146            let mma = ();
147            stream.expect_complete()?;
148            stream.expect_string(".ws")?;
149            let ws = ();
150            stream.expect_complete()?;
151            stream.expect_string(".cta_group::1")?;
152            let cta_group_1 = ();
153            stream.expect_complete()?;
154            let kind = Kind::parse(stream)?;
155            stream.expect_complete()?;
156            let saved_pos = stream.position();
157            let collector_usage = stream.expect_string(".collector_usage").is_ok();
158            if !collector_usage {
159                stream.set_position(saved_pos);
160            }
161            stream.expect_complete()?;
162            let d_tmem = AddressOperand::parse(stream)?;
163            stream.expect_complete()?;
164            stream.expect(&PtxToken::Comma)?;
165            let a_tmem = AddressOperand::parse(stream)?;
166            stream.expect_complete()?;
167            stream.expect(&PtxToken::Comma)?;
168            let b_desc = GeneralOperand::parse(stream)?;
169            stream.expect_complete()?;
170            stream.expect(&PtxToken::Comma)?;
171            let idesc = GeneralOperand::parse(stream)?;
172            stream.expect_complete()?;
173            stream.expect(&PtxToken::Comma)?;
174            let enable_input_d = GeneralOperand::parse(stream)?;
175            stream.expect_complete()?;
176            let saved_pos = stream.position();
177            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
178            if !has_comma {
179                stream.set_position(saved_pos);
180            }
181            let saved_pos = stream.position();
182            let zero_column_mask_desc = match GeneralOperand::parse(stream) {
183                Ok(val) => Some(val),
184                Err(_) => {
185                    stream.set_position(saved_pos);
186                    None
187                }
188            };
189            stream.expect_complete()?;
190            stream.expect_complete()?;
191            stream.expect(&PtxToken::Semicolon)?;
192            Ok(Tcgen05MmaWsCtaGroup1KindCollectorUsage1 {
193                mma,
194                ws,
195                cta_group_1,
196                kind,
197                collector_usage,
198                d_tmem,
199                a_tmem,
200                b_desc,
201                idesc,
202                enable_input_d,
203                zero_column_mask_desc,
204            })
205        }
206    }
207}
208
209pub mod section_1 {
210    use super::*;
211    use crate::r#type::instruction::tcgen05_mma_ws::section_1::*;
212
213    // ============================================================================
214    // Generated enum parsers
215    // ============================================================================
216
217    impl PtxParser for Buffer {
218        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
219            // Try B0
220            {
221                let saved_pos = stream.position();
222                if stream.expect_string("::b0").is_ok() {
223                    return Ok(Buffer::B0);
224                }
225                stream.set_position(saved_pos);
226            }
227            let saved_pos = stream.position();
228            // Try B1
229            {
230                let saved_pos = stream.position();
231                if stream.expect_string("::b1").is_ok() {
232                    return Ok(Buffer::B1);
233                }
234                stream.set_position(saved_pos);
235            }
236            stream.set_position(saved_pos);
237            let saved_pos = stream.position();
238            // Try B2
239            {
240                let saved_pos = stream.position();
241                if stream.expect_string("::b2").is_ok() {
242                    return Ok(Buffer::B2);
243                }
244                stream.set_position(saved_pos);
245            }
246            stream.set_position(saved_pos);
247            let saved_pos = stream.position();
248            // Try B3
249            {
250                let saved_pos = stream.position();
251                if stream.expect_string("::b3").is_ok() {
252                    return Ok(Buffer::B3);
253                }
254                stream.set_position(saved_pos);
255            }
256            stream.set_position(saved_pos);
257            let span = stream
258                .peek()
259                .map(|(_, s)| s.clone())
260                .unwrap_or(Span { start: 0, end: 0 });
261            let expected = &["::b0", "::b1", "::b2", "::b3"];
262            let found = stream
263                .peek()
264                .map(|(t, _)| format!("{:?}", t))
265                .unwrap_or_else(|_| "<end of input>".to_string());
266            Err(crate::parser::unexpected_value(span, expected, found))
267        }
268    }
269
270    impl PtxParser for CollectorUsage {
271        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
272            // Try CollectorBufferOp
273            {
274                let saved_seq_pos = stream.position();
275                match (|| -> Result<_, PtxParseError> {
276                    stream.expect_string(".collector")?;
277                    let collector = ();
278                    let buffer = Buffer::parse(stream)?;
279                    let op = Op::parse(stream)?;
280                    Ok((collector, buffer, op))
281                })() {
282                    Ok((collector, buffer, op)) => {
283                        return Ok(CollectorUsage::CollectorBufferOp(collector, buffer, op));
284                    }
285                    Err(_) => {
286                        stream.set_position(saved_seq_pos);
287                    }
288                }
289            }
290            let span = stream
291                .peek()
292                .map(|(_, s)| s.clone())
293                .unwrap_or(Span { start: 0, end: 0 });
294            let expected = &["<complex>"];
295            let found = stream
296                .peek()
297                .map(|(t, _)| format!("{:?}", t))
298                .unwrap_or_else(|_| "<end of input>".to_string());
299            Err(crate::parser::unexpected_value(span, expected, found))
300        }
301    }
302
303    impl PtxParser for Op {
304        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
305            // Try Lastuse
306            {
307                let saved_pos = stream.position();
308                if stream.expect_string("::lastuse").is_ok() {
309                    return Ok(Op::Lastuse);
310                }
311                stream.set_position(saved_pos);
312            }
313            let saved_pos = stream.position();
314            // Try Discard
315            {
316                let saved_pos = stream.position();
317                if stream.expect_string("::discard").is_ok() {
318                    return Ok(Op::Discard);
319                }
320                stream.set_position(saved_pos);
321            }
322            stream.set_position(saved_pos);
323            let saved_pos = stream.position();
324            // Try Fill
325            {
326                let saved_pos = stream.position();
327                if stream.expect_string("::fill").is_ok() {
328                    return Ok(Op::Fill);
329                }
330                stream.set_position(saved_pos);
331            }
332            stream.set_position(saved_pos);
333            let saved_pos = stream.position();
334            // Try Use
335            {
336                let saved_pos = stream.position();
337                if stream.expect_string("::use").is_ok() {
338                    return Ok(Op::Use);
339                }
340                stream.set_position(saved_pos);
341            }
342            stream.set_position(saved_pos);
343            let span = stream
344                .peek()
345                .map(|(_, s)| s.clone())
346                .unwrap_or(Span { start: 0, end: 0 });
347            let expected = &["::lastuse", "::discard", "::fill", "::use"];
348            let found = stream
349                .peek()
350                .map(|(t, _)| format!("{:?}", t))
351                .unwrap_or_else(|_| "<end of input>".to_string());
352            Err(crate::parser::unexpected_value(span, expected, found))
353        }
354    }
355
356    impl PtxParser for Tcgen05MmaWsCtaGroup1KindI8CollectorUsage {
357        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
358            stream.expect_string("tcgen05")?;
359            stream.expect_string(".mma")?;
360            let mma = ();
361            stream.expect_complete()?;
362            stream.expect_string(".ws")?;
363            let ws = ();
364            stream.expect_complete()?;
365            stream.expect_string(".cta_group::1")?;
366            let cta_group_1 = ();
367            stream.expect_complete()?;
368            stream.expect_string(".kind::i8")?;
369            let kind_i8 = ();
370            stream.expect_complete()?;
371            let saved_pos = stream.position();
372            let collector_usage = match CollectorUsage::parse(stream) {
373                Ok(val) => Some(val),
374                Err(_) => {
375                    stream.set_position(saved_pos);
376                    None
377                }
378            };
379            stream.expect_complete()?;
380            let d_tmem = AddressOperand::parse(stream)?;
381            stream.expect_complete()?;
382            stream.expect(&PtxToken::Comma)?;
383            let a_desc = GeneralOperand::parse(stream)?;
384            stream.expect_complete()?;
385            stream.expect(&PtxToken::Comma)?;
386            let b_desc = GeneralOperand::parse(stream)?;
387            stream.expect_complete()?;
388            stream.expect(&PtxToken::Comma)?;
389            let idesc = GeneralOperand::parse(stream)?;
390            stream.expect_complete()?;
391            stream.expect(&PtxToken::Comma)?;
392            let enable_input_d = GeneralOperand::parse(stream)?;
393            stream.expect_complete()?;
394            let saved_pos = stream.position();
395            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
396            if !has_comma {
397                stream.set_position(saved_pos);
398            }
399            let saved_pos = stream.position();
400            let zero_column_mask_desc = match GeneralOperand::parse(stream) {
401                Ok(val) => Some(val),
402                Err(_) => {
403                    stream.set_position(saved_pos);
404                    None
405                }
406            };
407            stream.expect_complete()?;
408            stream.expect_complete()?;
409            stream.expect(&PtxToken::Semicolon)?;
410            Ok(Tcgen05MmaWsCtaGroup1KindI8CollectorUsage {
411                mma,
412                ws,
413                cta_group_1,
414                kind_i8,
415                collector_usage,
416                d_tmem,
417                a_desc,
418                b_desc,
419                idesc,
420                enable_input_d,
421                zero_column_mask_desc,
422            })
423        }
424    }
425
426    impl PtxParser for Tcgen05MmaWsCtaGroup1KindI8CollectorUsage1 {
427        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
428            stream.expect_string("tcgen05")?;
429            stream.expect_string(".mma")?;
430            let mma = ();
431            stream.expect_complete()?;
432            stream.expect_string(".ws")?;
433            let ws = ();
434            stream.expect_complete()?;
435            stream.expect_string(".cta_group::1")?;
436            let cta_group_1 = ();
437            stream.expect_complete()?;
438            stream.expect_string(".kind::i8")?;
439            let kind_i8 = ();
440            stream.expect_complete()?;
441            let saved_pos = stream.position();
442            let collector_usage = match CollectorUsage::parse(stream) {
443                Ok(val) => Some(val),
444                Err(_) => {
445                    stream.set_position(saved_pos);
446                    None
447                }
448            };
449            stream.expect_complete()?;
450            let d_tmem = AddressOperand::parse(stream)?;
451            stream.expect_complete()?;
452            stream.expect(&PtxToken::Comma)?;
453            let a_tmem = AddressOperand::parse(stream)?;
454            stream.expect_complete()?;
455            stream.expect(&PtxToken::Comma)?;
456            let b_desc = GeneralOperand::parse(stream)?;
457            stream.expect_complete()?;
458            stream.expect(&PtxToken::Comma)?;
459            let idesc = GeneralOperand::parse(stream)?;
460            stream.expect_complete()?;
461            stream.expect(&PtxToken::Comma)?;
462            let enable_input_d = GeneralOperand::parse(stream)?;
463            stream.expect_complete()?;
464            let saved_pos = stream.position();
465            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
466            if !has_comma {
467                stream.set_position(saved_pos);
468            }
469            let saved_pos = stream.position();
470            let zero_column_mask_desc = match GeneralOperand::parse(stream) {
471                Ok(val) => Some(val),
472                Err(_) => {
473                    stream.set_position(saved_pos);
474                    None
475                }
476            };
477            stream.expect_complete()?;
478            stream.expect_complete()?;
479            stream.expect(&PtxToken::Semicolon)?;
480            Ok(Tcgen05MmaWsCtaGroup1KindI8CollectorUsage1 {
481                mma,
482                ws,
483                cta_group_1,
484                kind_i8,
485                collector_usage,
486                d_tmem,
487                a_tmem,
488                b_desc,
489                idesc,
490                enable_input_d,
491                zero_column_mask_desc,
492            })
493        }
494    }
495}
ptx_parser/parser/instruction/tcgen05_mma_ws.rs

ptx_parser/parser/instruction/
tcgen05_mma_ws.rs
