ptx_parser/parser/instruction/
tcgen05_mma.rs

1//! Original PTX specification:
2//!
3//! // 1. Floating-point type without block scaling:
4//! tcgen05.mma.cta_group.kind   [d-tmem],  a-desc,  b-desc, idesc
5//! {, disable-output-lane }, enable-input-d {, scale-input-d};
6//! tcgen05.mma.cta_group.kind   [d-tmem], [a-tmem], b-desc, idesc
7//! {, disable-output-lane }, enable-input-d {, scale-input-d};
8//! .kind      = { .kind::f16, .kind::tf32, .kind::f8f6f4 };
9//! .cta_group = { .cta_group::1, .cta_group::2 };
10//! ------------------------------------------------------------------
11//! // 2. Floating-point type with block scaling:
12//! tcgen05.mma.cta_group.kind.block_scale{.scale_vectorsize}
13//! [d-tmem],  a-desc,  b-desc, idesc,
14//! [scale-A-tmem], [scale-B-tmem], enable-input-d;
15//! tcgen05.mma.cta_group.kind.block_scale{.scale_vectorsize}
16//! [d-tmem], [a-tmem], b-desc, idesc,
17//! [scale-A-tmem], [scale-B-tmem], enable-input-d;
18//! .kind = { .kind::mxf8f6f4, .kind::mxf4, .kind::mxf4nvf4 };
19//! .cta_group      = { .cta_group::1,   .cta_group::2 };
20//! .scale_vectorsize = { .scale_vec::1X, .scale_vec::2X, .scale_vec::4X, .block16, .block32 };
21//! ------------------------------------------------------------------
22//! // 3. Convolution MMA for floating-point type without block scaling:
23//! tcgen05.mma.cta_group.kind.collector_usage [d-tmem],  a-desc,  b-desc, idesc
24//! {, disable-output-lane }, enable-input-d {, scale-input-d};
25//! tcgen05.mma.cta_group.kind{.ashift}.collector_usage [d-tmem], [a-tmem], b-desc, idesc
26//! {, disable-output-lane }, enable-input-d {, scale-input-d};
27//! tcgen05.mma.cta_group.kind.ashift{.collector_usage} [d-tmem], [a-tmem], b-desc, idesc
28//! {, disable-output-lane }, enable-input-d {, scale-input-d};
29//! .kind      = { .kind::f16, .kind::tf32, .kind::f8f6f4 };
30//! .cta_group = { .cta_group::1,   .cta_group::2 };
31//! .collector_usage = { .collector::buffer::op };
32//! ::buffer         = { ::a };
33//! ::op             = { ::fill, ::use, ::lastuse, ::discard* };
34//! ------------------------------------------------------------------
35//! // 4. Activation Stationary MMA for floating-point type with block scaling:
36//! tcgen05.mma.cta_group.kind.block_scale{.scale_vectorsize}.collector_usage
37//! [d-tmem],  a-desc,  b-desc, idesc,
38//! [scale-A-tmem], [scale-B-tmem], enable-input-d;
39//! tcgen05.mma.cta_group.kind.block_scale{.scale_vectorsize}.collector_usage
40//! [d-tmem], [a-tmem], b-desc, idesc,
41//! [scale-A-tmem], [scale-B-tmem], enable-input-d;
42//! .cta_group       = { .cta_group::1,   .cta_group::2 };
43//! .scale_vectorsize  = { .scale_vec::1X, .scale_vec::2X, .scale_vec::4X, .block16, .block32 };
44//! .kind            = { .kind::mxf8f6f4, .kind::mxf4, .kind::mxf4nvf4 };
45//! .collector_usage = { .collector::buffer::op };
46//! ::buffer         = { ::a };
47//! ::op             = { ::fill, ::use, ::lastuse, ::discard* };
48//! ------------------------------------------------------------------
49//! // 5. Integer type:
50//! tcgen05.mma.cta_group.kind::i8  [d-tmem],  a-desc,  b-desc, idesc
51//! {, disable-output-lane }, enable-input-d;
52//! tcgen05.mma.cta_group.kind::i8  [d-tmem], [a-tmem], b-desc, idesc
53//! {, disable-output-lane }, enable-input-d;
54//! .cta_group = { .cta_group::1,   .cta_group::2  };
55//! ------------------------------------------------------------------
56//! // 6. Convolution MMA for integer type:
57//! tcgen05.mma.cta_group.kind::i8.collector_usage          [d-tmem],  a-desc,  b-desc, idesc
58//! {, disable-output-lane }, enable-input-d;
59//! tcgen05.mma.cta_group.kind::i8.ashift{.collector_usage} [d-tmem], [a-tmem], b-desc, idesc
60//! {, disable-output-lane }, enable-input-d;
61//! tcgen05.mma.cta_group.kind::i8{.ashift}.collector_usage [d-tmem], [a-tmem], b-desc, idesc
62//! {, disable-output-lane }, enable-input-d;
63//! .cta_group       = { .cta_group::1,   .cta_group::2  };
64//! .collector_usage = { .collector::buffer::op };
65//! ::buffer         = { ::a };
66//! ::op             = { ::fill, ::use, ::lastuse, ::discard* };
67
68#![allow(unused)]
69
70use crate::lexer::PtxToken;
71use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
72use crate::r#type::common::*;
73
74pub mod section_0 {
75    use super::*;
76    use crate::r#type::instruction::tcgen05_mma::section_0::*;
77
78    // ============================================================================
79    // Generated enum parsers
80    // ============================================================================
81
82    impl PtxParser for CtaGroup {
83        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
84            // Try CtaGroup1
85            {
86                let saved_pos = stream.position();
87                if stream.expect_string(".cta_group::1").is_ok() {
88                    return Ok(CtaGroup::CtaGroup1);
89                }
90                stream.set_position(saved_pos);
91            }
92            let saved_pos = stream.position();
93            // Try CtaGroup2
94            {
95                let saved_pos = stream.position();
96                if stream.expect_string(".cta_group::2").is_ok() {
97                    return Ok(CtaGroup::CtaGroup2);
98                }
99                stream.set_position(saved_pos);
100            }
101            stream.set_position(saved_pos);
102            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
103            let expected = &[".cta_group::1", ".cta_group::2"];
104            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
105            Err(crate::parser::unexpected_value(span, expected, found))
106        }
107    }
108
109    impl PtxParser for Kind {
110        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
111            // Try KindF8f6f4
112            {
113                let saved_pos = stream.position();
114                if stream.expect_string(".kind::f8f6f4").is_ok() {
115                    return Ok(Kind::KindF8f6f4);
116                }
117                stream.set_position(saved_pos);
118            }
119            let saved_pos = stream.position();
120            // Try KindTf32
121            {
122                let saved_pos = stream.position();
123                if stream.expect_string(".kind::tf32").is_ok() {
124                    return Ok(Kind::KindTf32);
125                }
126                stream.set_position(saved_pos);
127            }
128            stream.set_position(saved_pos);
129            let saved_pos = stream.position();
130            // Try KindF16
131            {
132                let saved_pos = stream.position();
133                if stream.expect_string(".kind::f16").is_ok() {
134                    return Ok(Kind::KindF16);
135                }
136                stream.set_position(saved_pos);
137            }
138            stream.set_position(saved_pos);
139            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
140            let expected = &[".kind::f8f6f4", ".kind::tf32", ".kind::f16"];
141            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
142            Err(crate::parser::unexpected_value(span, expected, found))
143        }
144    }
145
146    impl PtxParser for Tcgen05MmaCtaGroupKind {
147        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
148            stream.expect_string("tcgen05")?;
149            stream.expect_string(".mma")?;
150            let mma = ();
151            stream.expect_complete()?;
152            let cta_group = CtaGroup::parse(stream)?;
153            stream.expect_complete()?;
154            let kind = Kind::parse(stream)?;
155            stream.expect_complete()?;
156            let d_tmem = AddressOperand::parse(stream)?;
157            stream.expect_complete()?;
158            stream.expect(&PtxToken::Comma)?;
159            let a_desc = GeneralOperand::parse(stream)?;
160            stream.expect_complete()?;
161            stream.expect(&PtxToken::Comma)?;
162            let b_desc = GeneralOperand::parse(stream)?;
163            stream.expect_complete()?;
164            stream.expect(&PtxToken::Comma)?;
165            let idesc = GeneralOperand::parse(stream)?;
166            stream.expect_complete()?;
167            let saved_pos = stream.position();
168            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
169            if !has_comma {
170                stream.set_position(saved_pos);
171            }
172            let saved_pos = stream.position();
173            let disable_output_lane = match GeneralOperand::parse(stream) {
174                Ok(val) => Some(val),
175                Err(_) => {
176                    stream.set_position(saved_pos);
177                    None
178                }
179            };
180            stream.expect_complete()?;
181            stream.expect(&PtxToken::Comma)?;
182            let enable_input_d = GeneralOperand::parse(stream)?;
183            stream.expect_complete()?;
184            let saved_pos = stream.position();
185            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
186            if !has_comma {
187                stream.set_position(saved_pos);
188            }
189            let saved_pos = stream.position();
190            let scale_input_d = match GeneralOperand::parse(stream) {
191                Ok(val) => Some(val),
192                Err(_) => {
193                    stream.set_position(saved_pos);
194                    None
195                }
196            };
197            stream.expect_complete()?;
198            stream.expect_complete()?;
199            stream.expect(&PtxToken::Semicolon)?;
200            Ok(Tcgen05MmaCtaGroupKind {
201                mma,
202                cta_group,
203                kind,
204                d_tmem,
205                a_desc,
206                b_desc,
207                idesc,
208                disable_output_lane,
209                enable_input_d,
210                scale_input_d,
211            })
212        }
213    }
214
215
216    impl PtxParser for Tcgen05MmaCtaGroupKind1 {
217        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
218            stream.expect_string("tcgen05")?;
219            stream.expect_string(".mma")?;
220            let mma = ();
221            stream.expect_complete()?;
222            let cta_group = CtaGroup::parse(stream)?;
223            stream.expect_complete()?;
224            let kind = Kind::parse(stream)?;
225            stream.expect_complete()?;
226            let d_tmem = AddressOperand::parse(stream)?;
227            stream.expect_complete()?;
228            stream.expect(&PtxToken::Comma)?;
229            let a_tmem = AddressOperand::parse(stream)?;
230            stream.expect_complete()?;
231            stream.expect(&PtxToken::Comma)?;
232            let b_desc = GeneralOperand::parse(stream)?;
233            stream.expect_complete()?;
234            stream.expect(&PtxToken::Comma)?;
235            let idesc = GeneralOperand::parse(stream)?;
236            stream.expect_complete()?;
237            let saved_pos = stream.position();
238            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
239            if !has_comma {
240                stream.set_position(saved_pos);
241            }
242            let saved_pos = stream.position();
243            let disable_output_lane = match GeneralOperand::parse(stream) {
244                Ok(val) => Some(val),
245                Err(_) => {
246                    stream.set_position(saved_pos);
247                    None
248                }
249            };
250            stream.expect_complete()?;
251            stream.expect(&PtxToken::Comma)?;
252            let enable_input_d = GeneralOperand::parse(stream)?;
253            stream.expect_complete()?;
254            let saved_pos = stream.position();
255            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
256            if !has_comma {
257                stream.set_position(saved_pos);
258            }
259            let saved_pos = stream.position();
260            let scale_input_d = match GeneralOperand::parse(stream) {
261                Ok(val) => Some(val),
262                Err(_) => {
263                    stream.set_position(saved_pos);
264                    None
265                }
266            };
267            stream.expect_complete()?;
268            stream.expect_complete()?;
269            stream.expect(&PtxToken::Semicolon)?;
270            Ok(Tcgen05MmaCtaGroupKind1 {
271                mma,
272                cta_group,
273                kind,
274                d_tmem,
275                a_tmem,
276                b_desc,
277                idesc,
278                disable_output_lane,
279                enable_input_d,
280                scale_input_d,
281            })
282        }
283    }
284
285
286}
287
288pub mod section_1 {
289    use super::*;
290    use crate::r#type::instruction::tcgen05_mma::section_1::*;
291
292    // ============================================================================
293    // Generated enum parsers
294    // ============================================================================
295
296    impl PtxParser for CtaGroup {
297        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
298            // Try CtaGroup1
299            {
300                let saved_pos = stream.position();
301                if stream.expect_string(".cta_group::1").is_ok() {
302                    return Ok(CtaGroup::CtaGroup1);
303                }
304                stream.set_position(saved_pos);
305            }
306            let saved_pos = stream.position();
307            // Try CtaGroup2
308            {
309                let saved_pos = stream.position();
310                if stream.expect_string(".cta_group::2").is_ok() {
311                    return Ok(CtaGroup::CtaGroup2);
312                }
313                stream.set_position(saved_pos);
314            }
315            stream.set_position(saved_pos);
316            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
317            let expected = &[".cta_group::1", ".cta_group::2"];
318            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
319            Err(crate::parser::unexpected_value(span, expected, found))
320        }
321    }
322
323    impl PtxParser for Kind {
324        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
325            // Try KindMxf8f6f4
326            {
327                let saved_pos = stream.position();
328                if stream.expect_string(".kind::mxf8f6f4").is_ok() {
329                    return Ok(Kind::KindMxf8f6f4);
330                }
331                stream.set_position(saved_pos);
332            }
333            let saved_pos = stream.position();
334            // Try KindMxf4nvf4
335            {
336                let saved_pos = stream.position();
337                if stream.expect_string(".kind::mxf4nvf4").is_ok() {
338                    return Ok(Kind::KindMxf4nvf4);
339                }
340                stream.set_position(saved_pos);
341            }
342            stream.set_position(saved_pos);
343            let saved_pos = stream.position();
344            // Try KindMxf4
345            {
346                let saved_pos = stream.position();
347                if stream.expect_string(".kind::mxf4").is_ok() {
348                    return Ok(Kind::KindMxf4);
349                }
350                stream.set_position(saved_pos);
351            }
352            stream.set_position(saved_pos);
353            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
354            let expected = &[".kind::mxf8f6f4", ".kind::mxf4nvf4", ".kind::mxf4"];
355            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
356            Err(crate::parser::unexpected_value(span, expected, found))
357        }
358    }
359
360    impl PtxParser for ScaleVectorsize {
361        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
362            // Try ScaleVec1x
363            {
364                let saved_pos = stream.position();
365                if stream.expect_string(".scale_vec::1X").is_ok() {
366                    return Ok(ScaleVectorsize::ScaleVec1x);
367                }
368                stream.set_position(saved_pos);
369            }
370            let saved_pos = stream.position();
371            // Try ScaleVec2x
372            {
373                let saved_pos = stream.position();
374                if stream.expect_string(".scale_vec::2X").is_ok() {
375                    return Ok(ScaleVectorsize::ScaleVec2x);
376                }
377                stream.set_position(saved_pos);
378            }
379            stream.set_position(saved_pos);
380            let saved_pos = stream.position();
381            // Try ScaleVec4x
382            {
383                let saved_pos = stream.position();
384                if stream.expect_string(".scale_vec::4X").is_ok() {
385                    return Ok(ScaleVectorsize::ScaleVec4x);
386                }
387                stream.set_position(saved_pos);
388            }
389            stream.set_position(saved_pos);
390            let saved_pos = stream.position();
391            // Try Block16
392            {
393                let saved_pos = stream.position();
394                if stream.expect_string(".block16").is_ok() {
395                    return Ok(ScaleVectorsize::Block16);
396                }
397                stream.set_position(saved_pos);
398            }
399            stream.set_position(saved_pos);
400            let saved_pos = stream.position();
401            // Try Block32
402            {
403                let saved_pos = stream.position();
404                if stream.expect_string(".block32").is_ok() {
405                    return Ok(ScaleVectorsize::Block32);
406                }
407                stream.set_position(saved_pos);
408            }
409            stream.set_position(saved_pos);
410            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
411            let expected = &[".scale_vec::1X", ".scale_vec::2X", ".scale_vec::4X", ".block16", ".block32"];
412            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
413            Err(crate::parser::unexpected_value(span, expected, found))
414        }
415    }
416
417    impl PtxParser for Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsize {
418        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
419            stream.expect_string("tcgen05")?;
420            stream.expect_string(".mma")?;
421            let mma = ();
422            stream.expect_complete()?;
423            let cta_group = CtaGroup::parse(stream)?;
424            stream.expect_complete()?;
425            let kind = Kind::parse(stream)?;
426            stream.expect_complete()?;
427            stream.expect_string(".block_scale")?;
428            let block_scale = ();
429            stream.expect_complete()?;
430            let saved_pos = stream.position();
431            let scale_vectorsize = match ScaleVectorsize::parse(stream) {
432                Ok(val) => Some(val),
433                Err(_) => {
434                    stream.set_position(saved_pos);
435                    None
436                }
437            };
438            stream.expect_complete()?;
439            let d_tmem = AddressOperand::parse(stream)?;
440            stream.expect_complete()?;
441            stream.expect(&PtxToken::Comma)?;
442            let a_desc = GeneralOperand::parse(stream)?;
443            stream.expect_complete()?;
444            stream.expect(&PtxToken::Comma)?;
445            let b_desc = GeneralOperand::parse(stream)?;
446            stream.expect_complete()?;
447            stream.expect(&PtxToken::Comma)?;
448            let idesc = GeneralOperand::parse(stream)?;
449            stream.expect_complete()?;
450            stream.expect(&PtxToken::Comma)?;
451            let scale_a_tmem = AddressOperand::parse(stream)?;
452            stream.expect_complete()?;
453            stream.expect(&PtxToken::Comma)?;
454            let scale_b_tmem = AddressOperand::parse(stream)?;
455            stream.expect_complete()?;
456            stream.expect(&PtxToken::Comma)?;
457            let enable_input_d = GeneralOperand::parse(stream)?;
458            stream.expect_complete()?;
459            stream.expect_complete()?;
460            stream.expect(&PtxToken::Semicolon)?;
461            Ok(Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsize {
462                mma,
463                cta_group,
464                kind,
465                block_scale,
466                scale_vectorsize,
467                d_tmem,
468                a_desc,
469                b_desc,
470                idesc,
471                scale_a_tmem,
472                scale_b_tmem,
473                enable_input_d,
474            })
475        }
476    }
477
478
479    impl PtxParser for Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsize1 {
480        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
481            stream.expect_string("tcgen05")?;
482            stream.expect_string(".mma")?;
483            let mma = ();
484            stream.expect_complete()?;
485            let cta_group = CtaGroup::parse(stream)?;
486            stream.expect_complete()?;
487            let kind = Kind::parse(stream)?;
488            stream.expect_complete()?;
489            stream.expect_string(".block_scale")?;
490            let block_scale = ();
491            stream.expect_complete()?;
492            let saved_pos = stream.position();
493            let scale_vectorsize = match ScaleVectorsize::parse(stream) {
494                Ok(val) => Some(val),
495                Err(_) => {
496                    stream.set_position(saved_pos);
497                    None
498                }
499            };
500            stream.expect_complete()?;
501            let d_tmem = AddressOperand::parse(stream)?;
502            stream.expect_complete()?;
503            stream.expect(&PtxToken::Comma)?;
504            let a_tmem = AddressOperand::parse(stream)?;
505            stream.expect_complete()?;
506            stream.expect(&PtxToken::Comma)?;
507            let b_desc = GeneralOperand::parse(stream)?;
508            stream.expect_complete()?;
509            stream.expect(&PtxToken::Comma)?;
510            let idesc = GeneralOperand::parse(stream)?;
511            stream.expect_complete()?;
512            stream.expect(&PtxToken::Comma)?;
513            let scale_a_tmem = AddressOperand::parse(stream)?;
514            stream.expect_complete()?;
515            stream.expect(&PtxToken::Comma)?;
516            let scale_b_tmem = AddressOperand::parse(stream)?;
517            stream.expect_complete()?;
518            stream.expect(&PtxToken::Comma)?;
519            let enable_input_d = GeneralOperand::parse(stream)?;
520            stream.expect_complete()?;
521            stream.expect_complete()?;
522            stream.expect(&PtxToken::Semicolon)?;
523            Ok(Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsize1 {
524                mma,
525                cta_group,
526                kind,
527                block_scale,
528                scale_vectorsize,
529                d_tmem,
530                a_tmem,
531                b_desc,
532                idesc,
533                scale_a_tmem,
534                scale_b_tmem,
535                enable_input_d,
536            })
537        }
538    }
539
540
541}
542
543pub mod section_2 {
544    use super::*;
545    use crate::r#type::instruction::tcgen05_mma::section_2::*;
546
547    // ============================================================================
548    // Generated enum parsers
549    // ============================================================================
550
551    impl PtxParser for Buffer {
552        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
553            // Try A
554            {
555                let saved_pos = stream.position();
556                if stream.expect_string("::a").is_ok() {
557                    return Ok(Buffer::A);
558                }
559                stream.set_position(saved_pos);
560            }
561            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
562            let expected = &["::a"];
563            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
564            Err(crate::parser::unexpected_value(span, expected, found))
565        }
566    }
567
568    impl PtxParser for CollectorUsage {
569        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
570            // Try CollectorBufferOp
571            {
572                let saved_seq_pos = stream.position();
573                match (|| -> Result<_, PtxParseError> {
574            stream.expect_string(".collector")?;
575            let collector = ();
576            let buffer = Buffer::parse(stream)?;
577            let op = Op::parse(stream)?;
578                    Ok((collector, buffer, op))
579                })() {
580                    Ok((collector, buffer, op)) => {
581                        return Ok(CollectorUsage::CollectorBufferOp(collector, buffer, op));
582                    }
583                    Err(_) => {
584                        stream.set_position(saved_seq_pos);
585                    }
586                }
587            }
588            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
589            let expected = &["<complex>"];
590            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
591            Err(crate::parser::unexpected_value(span, expected, found))
592        }
593    }
594
595    impl PtxParser for CtaGroup {
596        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
597            // Try CtaGroup1
598            {
599                let saved_pos = stream.position();
600                if stream.expect_string(".cta_group::1").is_ok() {
601                    return Ok(CtaGroup::CtaGroup1);
602                }
603                stream.set_position(saved_pos);
604            }
605            let saved_pos = stream.position();
606            // Try CtaGroup2
607            {
608                let saved_pos = stream.position();
609                if stream.expect_string(".cta_group::2").is_ok() {
610                    return Ok(CtaGroup::CtaGroup2);
611                }
612                stream.set_position(saved_pos);
613            }
614            stream.set_position(saved_pos);
615            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
616            let expected = &[".cta_group::1", ".cta_group::2"];
617            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
618            Err(crate::parser::unexpected_value(span, expected, found))
619        }
620    }
621
622    impl PtxParser for Kind {
623        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
624            // Try KindF8f6f4
625            {
626                let saved_pos = stream.position();
627                if stream.expect_string(".kind::f8f6f4").is_ok() {
628                    return Ok(Kind::KindF8f6f4);
629                }
630                stream.set_position(saved_pos);
631            }
632            let saved_pos = stream.position();
633            // Try KindTf32
634            {
635                let saved_pos = stream.position();
636                if stream.expect_string(".kind::tf32").is_ok() {
637                    return Ok(Kind::KindTf32);
638                }
639                stream.set_position(saved_pos);
640            }
641            stream.set_position(saved_pos);
642            let saved_pos = stream.position();
643            // Try KindF16
644            {
645                let saved_pos = stream.position();
646                if stream.expect_string(".kind::f16").is_ok() {
647                    return Ok(Kind::KindF16);
648                }
649                stream.set_position(saved_pos);
650            }
651            stream.set_position(saved_pos);
652            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
653            let expected = &[".kind::f8f6f4", ".kind::tf32", ".kind::f16"];
654            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
655            Err(crate::parser::unexpected_value(span, expected, found))
656        }
657    }
658
659    impl PtxParser for Op {
660        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
661            // Try Discard
662            {
663                let saved_pos = stream.position();
664                if stream.expect_string("::discard*").is_ok() {
665                    return Ok(Op::Discard);
666                }
667                stream.set_position(saved_pos);
668            }
669            let saved_pos = stream.position();
670            // Try Lastuse
671            {
672                let saved_pos = stream.position();
673                if stream.expect_string("::lastuse").is_ok() {
674                    return Ok(Op::Lastuse);
675                }
676                stream.set_position(saved_pos);
677            }
678            stream.set_position(saved_pos);
679            let saved_pos = stream.position();
680            // Try Fill
681            {
682                let saved_pos = stream.position();
683                if stream.expect_string("::fill").is_ok() {
684                    return Ok(Op::Fill);
685                }
686                stream.set_position(saved_pos);
687            }
688            stream.set_position(saved_pos);
689            let saved_pos = stream.position();
690            // Try Use
691            {
692                let saved_pos = stream.position();
693                if stream.expect_string("::use").is_ok() {
694                    return Ok(Op::Use);
695                }
696                stream.set_position(saved_pos);
697            }
698            stream.set_position(saved_pos);
699            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
700            let expected = &["::discard*", "::lastuse", "::fill", "::use"];
701            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
702            Err(crate::parser::unexpected_value(span, expected, found))
703        }
704    }
705
706    impl PtxParser for Tcgen05MmaCtaGroupKindCollectorUsage {
707        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
708            stream.expect_string("tcgen05")?;
709            stream.expect_string(".mma")?;
710            let mma = ();
711            stream.expect_complete()?;
712            let cta_group = CtaGroup::parse(stream)?;
713            stream.expect_complete()?;
714            let kind = Kind::parse(stream)?;
715            stream.expect_complete()?;
716            let collector_usage = CollectorUsage::parse(stream)?;
717            stream.expect_complete()?;
718            let d_tmem = AddressOperand::parse(stream)?;
719            stream.expect_complete()?;
720            stream.expect(&PtxToken::Comma)?;
721            let a_desc = GeneralOperand::parse(stream)?;
722            stream.expect_complete()?;
723            stream.expect(&PtxToken::Comma)?;
724            let b_desc = GeneralOperand::parse(stream)?;
725            stream.expect_complete()?;
726            stream.expect(&PtxToken::Comma)?;
727            let idesc = GeneralOperand::parse(stream)?;
728            stream.expect_complete()?;
729            let saved_pos = stream.position();
730            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
731            if !has_comma {
732                stream.set_position(saved_pos);
733            }
734            let saved_pos = stream.position();
735            let disable_output_lane = match GeneralOperand::parse(stream) {
736                Ok(val) => Some(val),
737                Err(_) => {
738                    stream.set_position(saved_pos);
739                    None
740                }
741            };
742            stream.expect_complete()?;
743            stream.expect(&PtxToken::Comma)?;
744            let enable_input_d = GeneralOperand::parse(stream)?;
745            stream.expect_complete()?;
746            let saved_pos = stream.position();
747            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
748            if !has_comma {
749                stream.set_position(saved_pos);
750            }
751            let saved_pos = stream.position();
752            let scale_input_d = match GeneralOperand::parse(stream) {
753                Ok(val) => Some(val),
754                Err(_) => {
755                    stream.set_position(saved_pos);
756                    None
757                }
758            };
759            stream.expect_complete()?;
760            stream.expect_complete()?;
761            stream.expect(&PtxToken::Semicolon)?;
762            Ok(Tcgen05MmaCtaGroupKindCollectorUsage {
763                mma,
764                cta_group,
765                kind,
766                collector_usage,
767                d_tmem,
768                a_desc,
769                b_desc,
770                idesc,
771                disable_output_lane,
772                enable_input_d,
773                scale_input_d,
774            })
775        }
776    }
777
778
779    impl PtxParser for Tcgen05MmaCtaGroupKindAshiftCollectorUsage {
780        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
781            stream.expect_string("tcgen05")?;
782            stream.expect_string(".mma")?;
783            let mma = ();
784            stream.expect_complete()?;
785            let cta_group = CtaGroup::parse(stream)?;
786            stream.expect_complete()?;
787            let kind = Kind::parse(stream)?;
788            stream.expect_complete()?;
789            let saved_pos = stream.position();
790            let ashift = stream.expect_string(".ashift").is_ok();
791            if !ashift {
792                stream.set_position(saved_pos);
793            }
794            stream.expect_complete()?;
795            let collector_usage = CollectorUsage::parse(stream)?;
796            stream.expect_complete()?;
797            let d_tmem = AddressOperand::parse(stream)?;
798            stream.expect_complete()?;
799            stream.expect(&PtxToken::Comma)?;
800            let a_tmem = AddressOperand::parse(stream)?;
801            stream.expect_complete()?;
802            stream.expect(&PtxToken::Comma)?;
803            let b_desc = GeneralOperand::parse(stream)?;
804            stream.expect_complete()?;
805            stream.expect(&PtxToken::Comma)?;
806            let idesc = GeneralOperand::parse(stream)?;
807            stream.expect_complete()?;
808            let saved_pos = stream.position();
809            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
810            if !has_comma {
811                stream.set_position(saved_pos);
812            }
813            let saved_pos = stream.position();
814            let disable_output_lane = match GeneralOperand::parse(stream) {
815                Ok(val) => Some(val),
816                Err(_) => {
817                    stream.set_position(saved_pos);
818                    None
819                }
820            };
821            stream.expect_complete()?;
822            stream.expect(&PtxToken::Comma)?;
823            let enable_input_d = GeneralOperand::parse(stream)?;
824            stream.expect_complete()?;
825            let saved_pos = stream.position();
826            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
827            if !has_comma {
828                stream.set_position(saved_pos);
829            }
830            let saved_pos = stream.position();
831            let scale_input_d = match GeneralOperand::parse(stream) {
832                Ok(val) => Some(val),
833                Err(_) => {
834                    stream.set_position(saved_pos);
835                    None
836                }
837            };
838            stream.expect_complete()?;
839            stream.expect_complete()?;
840            stream.expect(&PtxToken::Semicolon)?;
841            Ok(Tcgen05MmaCtaGroupKindAshiftCollectorUsage {
842                mma,
843                cta_group,
844                kind,
845                ashift,
846                collector_usage,
847                d_tmem,
848                a_tmem,
849                b_desc,
850                idesc,
851                disable_output_lane,
852                enable_input_d,
853                scale_input_d,
854            })
855        }
856    }
857
858
859    impl PtxParser for Tcgen05MmaCtaGroupKindAshiftCollectorUsage1 {
860        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
861            stream.expect_string("tcgen05")?;
862            stream.expect_string(".mma")?;
863            let mma = ();
864            stream.expect_complete()?;
865            let cta_group = CtaGroup::parse(stream)?;
866            stream.expect_complete()?;
867            let kind = Kind::parse(stream)?;
868            stream.expect_complete()?;
869            stream.expect_string(".ashift")?;
870            let ashift = ();
871            stream.expect_complete()?;
872            let saved_pos = stream.position();
873            let collector_usage = match CollectorUsage::parse(stream) {
874                Ok(val) => Some(val),
875                Err(_) => {
876                    stream.set_position(saved_pos);
877                    None
878                }
879            };
880            stream.expect_complete()?;
881            let d_tmem = AddressOperand::parse(stream)?;
882            stream.expect_complete()?;
883            stream.expect(&PtxToken::Comma)?;
884            let a_tmem = AddressOperand::parse(stream)?;
885            stream.expect_complete()?;
886            stream.expect(&PtxToken::Comma)?;
887            let b_desc = GeneralOperand::parse(stream)?;
888            stream.expect_complete()?;
889            stream.expect(&PtxToken::Comma)?;
890            let idesc = GeneralOperand::parse(stream)?;
891            stream.expect_complete()?;
892            let saved_pos = stream.position();
893            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
894            if !has_comma {
895                stream.set_position(saved_pos);
896            }
897            let saved_pos = stream.position();
898            let disable_output_lane = match GeneralOperand::parse(stream) {
899                Ok(val) => Some(val),
900                Err(_) => {
901                    stream.set_position(saved_pos);
902                    None
903                }
904            };
905            stream.expect_complete()?;
906            stream.expect(&PtxToken::Comma)?;
907            let enable_input_d = GeneralOperand::parse(stream)?;
908            stream.expect_complete()?;
909            let saved_pos = stream.position();
910            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
911            if !has_comma {
912                stream.set_position(saved_pos);
913            }
914            let saved_pos = stream.position();
915            let scale_input_d = match GeneralOperand::parse(stream) {
916                Ok(val) => Some(val),
917                Err(_) => {
918                    stream.set_position(saved_pos);
919                    None
920                }
921            };
922            stream.expect_complete()?;
923            stream.expect_complete()?;
924            stream.expect(&PtxToken::Semicolon)?;
925            Ok(Tcgen05MmaCtaGroupKindAshiftCollectorUsage1 {
926                mma,
927                cta_group,
928                kind,
929                ashift,
930                collector_usage,
931                d_tmem,
932                a_tmem,
933                b_desc,
934                idesc,
935                disable_output_lane,
936                enable_input_d,
937                scale_input_d,
938            })
939        }
940    }
941
942
943}
944
945pub mod section_3 {
946    use super::*;
947    use crate::r#type::instruction::tcgen05_mma::section_3::*;
948
949    // ============================================================================
950    // Generated enum parsers
951    // ============================================================================
952
953    impl PtxParser for Buffer {
954        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
955            // Try A
956            {
957                let saved_pos = stream.position();
958                if stream.expect_string("::a").is_ok() {
959                    return Ok(Buffer::A);
960                }
961                stream.set_position(saved_pos);
962            }
963            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
964            let expected = &["::a"];
965            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
966            Err(crate::parser::unexpected_value(span, expected, found))
967        }
968    }
969
970    impl PtxParser for CollectorUsage {
971        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
972            // Try CollectorBufferOp
973            {
974                let saved_seq_pos = stream.position();
975                match (|| -> Result<_, PtxParseError> {
976            stream.expect_string(".collector")?;
977            let collector = ();
978            let buffer = Buffer::parse(stream)?;
979            let op = Op::parse(stream)?;
980                    Ok((collector, buffer, op))
981                })() {
982                    Ok((collector, buffer, op)) => {
983                        return Ok(CollectorUsage::CollectorBufferOp(collector, buffer, op));
984                    }
985                    Err(_) => {
986                        stream.set_position(saved_seq_pos);
987                    }
988                }
989            }
990            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
991            let expected = &["<complex>"];
992            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
993            Err(crate::parser::unexpected_value(span, expected, found))
994        }
995    }
996
997    impl PtxParser for CtaGroup {
998        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
999            // Try CtaGroup1
1000            {
1001                let saved_pos = stream.position();
1002                if stream.expect_string(".cta_group::1").is_ok() {
1003                    return Ok(CtaGroup::CtaGroup1);
1004                }
1005                stream.set_position(saved_pos);
1006            }
1007            let saved_pos = stream.position();
1008            // Try CtaGroup2
1009            {
1010                let saved_pos = stream.position();
1011                if stream.expect_string(".cta_group::2").is_ok() {
1012                    return Ok(CtaGroup::CtaGroup2);
1013                }
1014                stream.set_position(saved_pos);
1015            }
1016            stream.set_position(saved_pos);
1017            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1018            let expected = &[".cta_group::1", ".cta_group::2"];
1019            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1020            Err(crate::parser::unexpected_value(span, expected, found))
1021        }
1022    }
1023
1024    impl PtxParser for Kind {
1025        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1026            // Try KindMxf8f6f4
1027            {
1028                let saved_pos = stream.position();
1029                if stream.expect_string(".kind::mxf8f6f4").is_ok() {
1030                    return Ok(Kind::KindMxf8f6f4);
1031                }
1032                stream.set_position(saved_pos);
1033            }
1034            let saved_pos = stream.position();
1035            // Try KindMxf4nvf4
1036            {
1037                let saved_pos = stream.position();
1038                if stream.expect_string(".kind::mxf4nvf4").is_ok() {
1039                    return Ok(Kind::KindMxf4nvf4);
1040                }
1041                stream.set_position(saved_pos);
1042            }
1043            stream.set_position(saved_pos);
1044            let saved_pos = stream.position();
1045            // Try KindMxf4
1046            {
1047                let saved_pos = stream.position();
1048                if stream.expect_string(".kind::mxf4").is_ok() {
1049                    return Ok(Kind::KindMxf4);
1050                }
1051                stream.set_position(saved_pos);
1052            }
1053            stream.set_position(saved_pos);
1054            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1055            let expected = &[".kind::mxf8f6f4", ".kind::mxf4nvf4", ".kind::mxf4"];
1056            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1057            Err(crate::parser::unexpected_value(span, expected, found))
1058        }
1059    }
1060
1061    impl PtxParser for Op {
1062        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1063            // Try Discard
1064            {
1065                let saved_pos = stream.position();
1066                if stream.expect_string("::discard*").is_ok() {
1067                    return Ok(Op::Discard);
1068                }
1069                stream.set_position(saved_pos);
1070            }
1071            let saved_pos = stream.position();
1072            // Try Lastuse
1073            {
1074                let saved_pos = stream.position();
1075                if stream.expect_string("::lastuse").is_ok() {
1076                    return Ok(Op::Lastuse);
1077                }
1078                stream.set_position(saved_pos);
1079            }
1080            stream.set_position(saved_pos);
1081            let saved_pos = stream.position();
1082            // Try Fill
1083            {
1084                let saved_pos = stream.position();
1085                if stream.expect_string("::fill").is_ok() {
1086                    return Ok(Op::Fill);
1087                }
1088                stream.set_position(saved_pos);
1089            }
1090            stream.set_position(saved_pos);
1091            let saved_pos = stream.position();
1092            // Try Use
1093            {
1094                let saved_pos = stream.position();
1095                if stream.expect_string("::use").is_ok() {
1096                    return Ok(Op::Use);
1097                }
1098                stream.set_position(saved_pos);
1099            }
1100            stream.set_position(saved_pos);
1101            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1102            let expected = &["::discard*", "::lastuse", "::fill", "::use"];
1103            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1104            Err(crate::parser::unexpected_value(span, expected, found))
1105        }
1106    }
1107
1108    impl PtxParser for ScaleVectorsize {
1109        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1110            // Try ScaleVec1x
1111            {
1112                let saved_pos = stream.position();
1113                if stream.expect_string(".scale_vec::1X").is_ok() {
1114                    return Ok(ScaleVectorsize::ScaleVec1x);
1115                }
1116                stream.set_position(saved_pos);
1117            }
1118            let saved_pos = stream.position();
1119            // Try ScaleVec2x
1120            {
1121                let saved_pos = stream.position();
1122                if stream.expect_string(".scale_vec::2X").is_ok() {
1123                    return Ok(ScaleVectorsize::ScaleVec2x);
1124                }
1125                stream.set_position(saved_pos);
1126            }
1127            stream.set_position(saved_pos);
1128            let saved_pos = stream.position();
1129            // Try ScaleVec4x
1130            {
1131                let saved_pos = stream.position();
1132                if stream.expect_string(".scale_vec::4X").is_ok() {
1133                    return Ok(ScaleVectorsize::ScaleVec4x);
1134                }
1135                stream.set_position(saved_pos);
1136            }
1137            stream.set_position(saved_pos);
1138            let saved_pos = stream.position();
1139            // Try Block16
1140            {
1141                let saved_pos = stream.position();
1142                if stream.expect_string(".block16").is_ok() {
1143                    return Ok(ScaleVectorsize::Block16);
1144                }
1145                stream.set_position(saved_pos);
1146            }
1147            stream.set_position(saved_pos);
1148            let saved_pos = stream.position();
1149            // Try Block32
1150            {
1151                let saved_pos = stream.position();
1152                if stream.expect_string(".block32").is_ok() {
1153                    return Ok(ScaleVectorsize::Block32);
1154                }
1155                stream.set_position(saved_pos);
1156            }
1157            stream.set_position(saved_pos);
1158            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1159            let expected = &[".scale_vec::1X", ".scale_vec::2X", ".scale_vec::4X", ".block16", ".block32"];
1160            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1161            Err(crate::parser::unexpected_value(span, expected, found))
1162        }
1163    }
1164
1165    impl PtxParser for Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage {
1166        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1167            stream.expect_string("tcgen05")?;
1168            stream.expect_string(".mma")?;
1169            let mma = ();
1170            stream.expect_complete()?;
1171            let cta_group = CtaGroup::parse(stream)?;
1172            stream.expect_complete()?;
1173            let kind = Kind::parse(stream)?;
1174            stream.expect_complete()?;
1175            stream.expect_string(".block_scale")?;
1176            let block_scale = ();
1177            stream.expect_complete()?;
1178            let saved_pos = stream.position();
1179            let scale_vectorsize = match ScaleVectorsize::parse(stream) {
1180                Ok(val) => Some(val),
1181                Err(_) => {
1182                    stream.set_position(saved_pos);
1183                    None
1184                }
1185            };
1186            stream.expect_complete()?;
1187            let collector_usage = CollectorUsage::parse(stream)?;
1188            stream.expect_complete()?;
1189            let d_tmem = AddressOperand::parse(stream)?;
1190            stream.expect_complete()?;
1191            stream.expect(&PtxToken::Comma)?;
1192            let a_desc = GeneralOperand::parse(stream)?;
1193            stream.expect_complete()?;
1194            stream.expect(&PtxToken::Comma)?;
1195            let b_desc = GeneralOperand::parse(stream)?;
1196            stream.expect_complete()?;
1197            stream.expect(&PtxToken::Comma)?;
1198            let idesc = GeneralOperand::parse(stream)?;
1199            stream.expect_complete()?;
1200            stream.expect(&PtxToken::Comma)?;
1201            let scale_a_tmem = AddressOperand::parse(stream)?;
1202            stream.expect_complete()?;
1203            stream.expect(&PtxToken::Comma)?;
1204            let scale_b_tmem = AddressOperand::parse(stream)?;
1205            stream.expect_complete()?;
1206            stream.expect(&PtxToken::Comma)?;
1207            let enable_input_d = GeneralOperand::parse(stream)?;
1208            stream.expect_complete()?;
1209            stream.expect_complete()?;
1210            stream.expect(&PtxToken::Semicolon)?;
1211            Ok(Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage {
1212                mma,
1213                cta_group,
1214                kind,
1215                block_scale,
1216                scale_vectorsize,
1217                collector_usage,
1218                d_tmem,
1219                a_desc,
1220                b_desc,
1221                idesc,
1222                scale_a_tmem,
1223                scale_b_tmem,
1224                enable_input_d,
1225            })
1226        }
1227    }
1228
1229
1230    impl PtxParser for Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage1 {
1231        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1232            stream.expect_string("tcgen05")?;
1233            stream.expect_string(".mma")?;
1234            let mma = ();
1235            stream.expect_complete()?;
1236            let cta_group = CtaGroup::parse(stream)?;
1237            stream.expect_complete()?;
1238            let kind = Kind::parse(stream)?;
1239            stream.expect_complete()?;
1240            stream.expect_string(".block_scale")?;
1241            let block_scale = ();
1242            stream.expect_complete()?;
1243            let saved_pos = stream.position();
1244            let scale_vectorsize = match ScaleVectorsize::parse(stream) {
1245                Ok(val) => Some(val),
1246                Err(_) => {
1247                    stream.set_position(saved_pos);
1248                    None
1249                }
1250            };
1251            stream.expect_complete()?;
1252            let collector_usage = CollectorUsage::parse(stream)?;
1253            stream.expect_complete()?;
1254            let d_tmem = AddressOperand::parse(stream)?;
1255            stream.expect_complete()?;
1256            stream.expect(&PtxToken::Comma)?;
1257            let a_tmem = AddressOperand::parse(stream)?;
1258            stream.expect_complete()?;
1259            stream.expect(&PtxToken::Comma)?;
1260            let b_desc = GeneralOperand::parse(stream)?;
1261            stream.expect_complete()?;
1262            stream.expect(&PtxToken::Comma)?;
1263            let idesc = GeneralOperand::parse(stream)?;
1264            stream.expect_complete()?;
1265            stream.expect(&PtxToken::Comma)?;
1266            let scale_a_tmem = AddressOperand::parse(stream)?;
1267            stream.expect_complete()?;
1268            stream.expect(&PtxToken::Comma)?;
1269            let scale_b_tmem = AddressOperand::parse(stream)?;
1270            stream.expect_complete()?;
1271            stream.expect(&PtxToken::Comma)?;
1272            let enable_input_d = GeneralOperand::parse(stream)?;
1273            stream.expect_complete()?;
1274            stream.expect_complete()?;
1275            stream.expect(&PtxToken::Semicolon)?;
1276            Ok(Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage1 {
1277                mma,
1278                cta_group,
1279                kind,
1280                block_scale,
1281                scale_vectorsize,
1282                collector_usage,
1283                d_tmem,
1284                a_tmem,
1285                b_desc,
1286                idesc,
1287                scale_a_tmem,
1288                scale_b_tmem,
1289                enable_input_d,
1290            })
1291        }
1292    }
1293
1294
1295}
1296
1297pub mod section_4 {
1298    use super::*;
1299    use crate::r#type::instruction::tcgen05_mma::section_4::*;
1300
1301    // ============================================================================
1302    // Generated enum parsers
1303    // ============================================================================
1304
1305    impl PtxParser for CtaGroup {
1306        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1307            // Try CtaGroup1
1308            {
1309                let saved_pos = stream.position();
1310                if stream.expect_string(".cta_group::1").is_ok() {
1311                    return Ok(CtaGroup::CtaGroup1);
1312                }
1313                stream.set_position(saved_pos);
1314            }
1315            let saved_pos = stream.position();
1316            // Try CtaGroup2
1317            {
1318                let saved_pos = stream.position();
1319                if stream.expect_string(".cta_group::2").is_ok() {
1320                    return Ok(CtaGroup::CtaGroup2);
1321                }
1322                stream.set_position(saved_pos);
1323            }
1324            stream.set_position(saved_pos);
1325            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1326            let expected = &[".cta_group::1", ".cta_group::2"];
1327            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1328            Err(crate::parser::unexpected_value(span, expected, found))
1329        }
1330    }
1331
1332    impl PtxParser for Tcgen05MmaCtaGroupKindI8 {
1333        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1334            stream.expect_string("tcgen05")?;
1335            stream.expect_string(".mma")?;
1336            let mma = ();
1337            stream.expect_complete()?;
1338            let cta_group = CtaGroup::parse(stream)?;
1339            stream.expect_complete()?;
1340            stream.expect_string(".kind::i8")?;
1341            let kind_i8 = ();
1342            stream.expect_complete()?;
1343            let d_tmem = AddressOperand::parse(stream)?;
1344            stream.expect_complete()?;
1345            stream.expect(&PtxToken::Comma)?;
1346            let a_desc = GeneralOperand::parse(stream)?;
1347            stream.expect_complete()?;
1348            stream.expect(&PtxToken::Comma)?;
1349            let b_desc = GeneralOperand::parse(stream)?;
1350            stream.expect_complete()?;
1351            stream.expect(&PtxToken::Comma)?;
1352            let idesc = GeneralOperand::parse(stream)?;
1353            stream.expect_complete()?;
1354            let saved_pos = stream.position();
1355            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
1356            if !has_comma {
1357                stream.set_position(saved_pos);
1358            }
1359            let saved_pos = stream.position();
1360            let disable_output_lane = match GeneralOperand::parse(stream) {
1361                Ok(val) => Some(val),
1362                Err(_) => {
1363                    stream.set_position(saved_pos);
1364                    None
1365                }
1366            };
1367            stream.expect_complete()?;
1368            stream.expect(&PtxToken::Comma)?;
1369            let enable_input_d = GeneralOperand::parse(stream)?;
1370            stream.expect_complete()?;
1371            stream.expect_complete()?;
1372            stream.expect(&PtxToken::Semicolon)?;
1373            Ok(Tcgen05MmaCtaGroupKindI8 {
1374                mma,
1375                cta_group,
1376                kind_i8,
1377                d_tmem,
1378                a_desc,
1379                b_desc,
1380                idesc,
1381                disable_output_lane,
1382                enable_input_d,
1383            })
1384        }
1385    }
1386
1387
1388    impl PtxParser for Tcgen05MmaCtaGroupKindI81 {
1389        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1390            stream.expect_string("tcgen05")?;
1391            stream.expect_string(".mma")?;
1392            let mma = ();
1393            stream.expect_complete()?;
1394            let cta_group = CtaGroup::parse(stream)?;
1395            stream.expect_complete()?;
1396            stream.expect_string(".kind::i8")?;
1397            let kind_i8 = ();
1398            stream.expect_complete()?;
1399            let d_tmem = AddressOperand::parse(stream)?;
1400            stream.expect_complete()?;
1401            stream.expect(&PtxToken::Comma)?;
1402            let a_tmem = AddressOperand::parse(stream)?;
1403            stream.expect_complete()?;
1404            stream.expect(&PtxToken::Comma)?;
1405            let b_desc = GeneralOperand::parse(stream)?;
1406            stream.expect_complete()?;
1407            stream.expect(&PtxToken::Comma)?;
1408            let idesc = GeneralOperand::parse(stream)?;
1409            stream.expect_complete()?;
1410            let saved_pos = stream.position();
1411            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
1412            if !has_comma {
1413                stream.set_position(saved_pos);
1414            }
1415            let saved_pos = stream.position();
1416            let disable_output_lane = match GeneralOperand::parse(stream) {
1417                Ok(val) => Some(val),
1418                Err(_) => {
1419                    stream.set_position(saved_pos);
1420                    None
1421                }
1422            };
1423            stream.expect_complete()?;
1424            stream.expect(&PtxToken::Comma)?;
1425            let enable_input_d = GeneralOperand::parse(stream)?;
1426            stream.expect_complete()?;
1427            stream.expect_complete()?;
1428            stream.expect(&PtxToken::Semicolon)?;
1429            Ok(Tcgen05MmaCtaGroupKindI81 {
1430                mma,
1431                cta_group,
1432                kind_i8,
1433                d_tmem,
1434                a_tmem,
1435                b_desc,
1436                idesc,
1437                disable_output_lane,
1438                enable_input_d,
1439            })
1440        }
1441    }
1442
1443
1444}
1445
1446pub mod section_5 {
1447    use super::*;
1448    use crate::r#type::instruction::tcgen05_mma::section_5::*;
1449
1450    // ============================================================================
1451    // Generated enum parsers
1452    // ============================================================================
1453
1454    impl PtxParser for Buffer {
1455        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1456            // Try A
1457            {
1458                let saved_pos = stream.position();
1459                if stream.expect_string("::a").is_ok() {
1460                    return Ok(Buffer::A);
1461                }
1462                stream.set_position(saved_pos);
1463            }
1464            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1465            let expected = &["::a"];
1466            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1467            Err(crate::parser::unexpected_value(span, expected, found))
1468        }
1469    }
1470
1471    impl PtxParser for CollectorUsage {
1472        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1473            // Try CollectorBufferOp
1474            {
1475                let saved_seq_pos = stream.position();
1476                match (|| -> Result<_, PtxParseError> {
1477            stream.expect_string(".collector")?;
1478            let collector = ();
1479            let buffer = Buffer::parse(stream)?;
1480            let op = Op::parse(stream)?;
1481                    Ok((collector, buffer, op))
1482                })() {
1483                    Ok((collector, buffer, op)) => {
1484                        return Ok(CollectorUsage::CollectorBufferOp(collector, buffer, op));
1485                    }
1486                    Err(_) => {
1487                        stream.set_position(saved_seq_pos);
1488                    }
1489                }
1490            }
1491            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1492            let expected = &["<complex>"];
1493            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1494            Err(crate::parser::unexpected_value(span, expected, found))
1495        }
1496    }
1497
1498    impl PtxParser for CtaGroup {
1499        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1500            // Try CtaGroup1
1501            {
1502                let saved_pos = stream.position();
1503                if stream.expect_string(".cta_group::1").is_ok() {
1504                    return Ok(CtaGroup::CtaGroup1);
1505                }
1506                stream.set_position(saved_pos);
1507            }
1508            let saved_pos = stream.position();
1509            // Try CtaGroup2
1510            {
1511                let saved_pos = stream.position();
1512                if stream.expect_string(".cta_group::2").is_ok() {
1513                    return Ok(CtaGroup::CtaGroup2);
1514                }
1515                stream.set_position(saved_pos);
1516            }
1517            stream.set_position(saved_pos);
1518            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1519            let expected = &[".cta_group::1", ".cta_group::2"];
1520            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1521            Err(crate::parser::unexpected_value(span, expected, found))
1522        }
1523    }
1524
1525    impl PtxParser for Op {
1526        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1527            // Try Discard
1528            {
1529                let saved_pos = stream.position();
1530                if stream.expect_string("::discard*").is_ok() {
1531                    return Ok(Op::Discard);
1532                }
1533                stream.set_position(saved_pos);
1534            }
1535            let saved_pos = stream.position();
1536            // Try Lastuse
1537            {
1538                let saved_pos = stream.position();
1539                if stream.expect_string("::lastuse").is_ok() {
1540                    return Ok(Op::Lastuse);
1541                }
1542                stream.set_position(saved_pos);
1543            }
1544            stream.set_position(saved_pos);
1545            let saved_pos = stream.position();
1546            // Try Fill
1547            {
1548                let saved_pos = stream.position();
1549                if stream.expect_string("::fill").is_ok() {
1550                    return Ok(Op::Fill);
1551                }
1552                stream.set_position(saved_pos);
1553            }
1554            stream.set_position(saved_pos);
1555            let saved_pos = stream.position();
1556            // Try Use
1557            {
1558                let saved_pos = stream.position();
1559                if stream.expect_string("::use").is_ok() {
1560                    return Ok(Op::Use);
1561                }
1562                stream.set_position(saved_pos);
1563            }
1564            stream.set_position(saved_pos);
1565            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1566            let expected = &["::discard*", "::lastuse", "::fill", "::use"];
1567            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1568            Err(crate::parser::unexpected_value(span, expected, found))
1569        }
1570    }
1571
1572    impl PtxParser for Tcgen05MmaCtaGroupKindI8CollectorUsage {
1573        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1574            stream.expect_string("tcgen05")?;
1575            stream.expect_string(".mma")?;
1576            let mma = ();
1577            stream.expect_complete()?;
1578            let cta_group = CtaGroup::parse(stream)?;
1579            stream.expect_complete()?;
1580            stream.expect_string(".kind::i8")?;
1581            let kind_i8 = ();
1582            stream.expect_complete()?;
1583            let collector_usage = CollectorUsage::parse(stream)?;
1584            stream.expect_complete()?;
1585            let d_tmem = AddressOperand::parse(stream)?;
1586            stream.expect_complete()?;
1587            stream.expect(&PtxToken::Comma)?;
1588            let a_desc = GeneralOperand::parse(stream)?;
1589            stream.expect_complete()?;
1590            stream.expect(&PtxToken::Comma)?;
1591            let b_desc = GeneralOperand::parse(stream)?;
1592            stream.expect_complete()?;
1593            stream.expect(&PtxToken::Comma)?;
1594            let idesc = GeneralOperand::parse(stream)?;
1595            stream.expect_complete()?;
1596            let saved_pos = stream.position();
1597            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
1598            if !has_comma {
1599                stream.set_position(saved_pos);
1600            }
1601            let saved_pos = stream.position();
1602            let disable_output_lane = match GeneralOperand::parse(stream) {
1603                Ok(val) => Some(val),
1604                Err(_) => {
1605                    stream.set_position(saved_pos);
1606                    None
1607                }
1608            };
1609            stream.expect_complete()?;
1610            stream.expect(&PtxToken::Comma)?;
1611            let enable_input_d = GeneralOperand::parse(stream)?;
1612            stream.expect_complete()?;
1613            stream.expect_complete()?;
1614            stream.expect(&PtxToken::Semicolon)?;
1615            Ok(Tcgen05MmaCtaGroupKindI8CollectorUsage {
1616                mma,
1617                cta_group,
1618                kind_i8,
1619                collector_usage,
1620                d_tmem,
1621                a_desc,
1622                b_desc,
1623                idesc,
1624                disable_output_lane,
1625                enable_input_d,
1626            })
1627        }
1628    }
1629
1630
1631    impl PtxParser for Tcgen05MmaCtaGroupKindI8AshiftCollectorUsage {
1632        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1633            stream.expect_string("tcgen05")?;
1634            stream.expect_string(".mma")?;
1635            let mma = ();
1636            stream.expect_complete()?;
1637            let cta_group = CtaGroup::parse(stream)?;
1638            stream.expect_complete()?;
1639            stream.expect_string(".kind::i8")?;
1640            let kind_i8 = ();
1641            stream.expect_complete()?;
1642            stream.expect_string(".ashift")?;
1643            let ashift = ();
1644            stream.expect_complete()?;
1645            let saved_pos = stream.position();
1646            let collector_usage = match CollectorUsage::parse(stream) {
1647                Ok(val) => Some(val),
1648                Err(_) => {
1649                    stream.set_position(saved_pos);
1650                    None
1651                }
1652            };
1653            stream.expect_complete()?;
1654            let d_tmem = AddressOperand::parse(stream)?;
1655            stream.expect_complete()?;
1656            stream.expect(&PtxToken::Comma)?;
1657            let a_tmem = AddressOperand::parse(stream)?;
1658            stream.expect_complete()?;
1659            stream.expect(&PtxToken::Comma)?;
1660            let b_desc = GeneralOperand::parse(stream)?;
1661            stream.expect_complete()?;
1662            stream.expect(&PtxToken::Comma)?;
1663            let idesc = GeneralOperand::parse(stream)?;
1664            stream.expect_complete()?;
1665            let saved_pos = stream.position();
1666            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
1667            if !has_comma {
1668                stream.set_position(saved_pos);
1669            }
1670            let saved_pos = stream.position();
1671            let disable_output_lane = match GeneralOperand::parse(stream) {
1672                Ok(val) => Some(val),
1673                Err(_) => {
1674                    stream.set_position(saved_pos);
1675                    None
1676                }
1677            };
1678            stream.expect_complete()?;
1679            stream.expect(&PtxToken::Comma)?;
1680            let enable_input_d = GeneralOperand::parse(stream)?;
1681            stream.expect_complete()?;
1682            stream.expect_complete()?;
1683            stream.expect(&PtxToken::Semicolon)?;
1684            Ok(Tcgen05MmaCtaGroupKindI8AshiftCollectorUsage {
1685                mma,
1686                cta_group,
1687                kind_i8,
1688                ashift,
1689                collector_usage,
1690                d_tmem,
1691                a_tmem,
1692                b_desc,
1693                idesc,
1694                disable_output_lane,
1695                enable_input_d,
1696            })
1697        }
1698    }
1699
1700
1701    impl PtxParser for Tcgen05MmaCtaGroupKindI8AshiftCollectorUsage1 {
1702        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1703            stream.expect_string("tcgen05")?;
1704            stream.expect_string(".mma")?;
1705            let mma = ();
1706            stream.expect_complete()?;
1707            let cta_group = CtaGroup::parse(stream)?;
1708            stream.expect_complete()?;
1709            stream.expect_string(".kind::i8")?;
1710            let kind_i8 = ();
1711            stream.expect_complete()?;
1712            let saved_pos = stream.position();
1713            let ashift = stream.expect_string(".ashift").is_ok();
1714            if !ashift {
1715                stream.set_position(saved_pos);
1716            }
1717            stream.expect_complete()?;
1718            let collector_usage = CollectorUsage::parse(stream)?;
1719            stream.expect_complete()?;
1720            let d_tmem = AddressOperand::parse(stream)?;
1721            stream.expect_complete()?;
1722            stream.expect(&PtxToken::Comma)?;
1723            let a_tmem = AddressOperand::parse(stream)?;
1724            stream.expect_complete()?;
1725            stream.expect(&PtxToken::Comma)?;
1726            let b_desc = GeneralOperand::parse(stream)?;
1727            stream.expect_complete()?;
1728            stream.expect(&PtxToken::Comma)?;
1729            let idesc = GeneralOperand::parse(stream)?;
1730            stream.expect_complete()?;
1731            let saved_pos = stream.position();
1732            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
1733            if !has_comma {
1734                stream.set_position(saved_pos);
1735            }
1736            let saved_pos = stream.position();
1737            let disable_output_lane = match GeneralOperand::parse(stream) {
1738                Ok(val) => Some(val),
1739                Err(_) => {
1740                    stream.set_position(saved_pos);
1741                    None
1742                }
1743            };
1744            stream.expect_complete()?;
1745            stream.expect(&PtxToken::Comma)?;
1746            let enable_input_d = GeneralOperand::parse(stream)?;
1747            stream.expect_complete()?;
1748            stream.expect_complete()?;
1749            stream.expect(&PtxToken::Semicolon)?;
1750            Ok(Tcgen05MmaCtaGroupKindI8AshiftCollectorUsage1 {
1751                mma,
1752                cta_group,
1753                kind_i8,
1754                ashift,
1755                collector_usage,
1756                d_tmem,
1757                a_tmem,
1758                b_desc,
1759                idesc,
1760                disable_output_lane,
1761                enable_input_d,
1762            })
1763        }
1764    }
1765
1766
1767}
1768
ptx_parser/parser/instruction/tcgen05_mma.rs

ptx_parser/parser/instruction/
tcgen05_mma.rs
