ptx_parser/parser/instruction/
tcgen05_mma_sp.rs

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

ptx_parser/parser/instruction/
tcgen05_mma_sp.rs
