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

ptx_parser/parser/instruction/
tcgen05_mma_ws_sp.rs
