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

ptx_parser/parser/instruction/
tcgen05_mma.rs
