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

ptx_parser/parser/instruction/
tcgen05_mma_sp.rs
