ptx-90-parser 0.3.0

//! Original PTX specification:
//!
//! // Half precision floating point type:
//! wgmma.mma_async.sp.sync.aligned.shape.dtype.f16.f16  d, a-desc, b-desc, sp-meta, sp-sel, scale-d, imm-scale-a, imm-scale-b, imm-trans-a, imm-trans-b;
//! wgmma.mma_async.sp.sync.aligned.shape.dtype.f16.f16  d, a, b-desc, sp-meta, sp-sel, scale-d, imm-scale-a, imm-scale-b, imm-trans-b;
//! .shape   = {.m64n8k32, .m64n16k32, .m64n24k32, .m64n32k32,
//! .m64n40k32, .m64n48k32, .m64n56k32, .m64n64k32,
//! .m64n72k32, .m64n80k32, .m64n88k32, .m64n96k32,
//! .m64n104k32, .m64n112k32, .m64n120k32, .m64n128k32,
//! .m64n136k32, .m64n144k32, .m64n152k32, .m64n160k32,
//! .m64n168k32, .m64n176k32, .m64n184k32, .m64n192k32,
//! .m64n200k32, .m64n208k32, .m64n216k32, .m64n224k32,
//! .m64n232k32, .m64n240k32, .m64n248k32, .m64n256k32};
//! .dtype   = {.f16, .f32};
//! ------------------------------------------------------------------
//! // Alternate floating point type :
//! // .bf16 floating point type:
//! wgmma.mma_async.sp.sync.aligned.shape.dtype.bf16.bf16  d, a-desc, b-desc, sp-meta, sp-sel, scale-d, imm-scale-a, imm-scale-b, imm-trans-a, imm-trans-b;
//! wgmma.mma_async.sp.sync.aligned.shape.dtype.bf16.bf16  d, a, b-desc, sp-meta, sp-sel, scale-d, imm-scale-a, imm-scale-b, imm-trans-b;
//! .shape   = {.m64n8k32, .m64n16k32, .m64n24k32, .m64n32k32,
//! .m64n40k32, .m64n48k32, .m64n56k32, .m64n64k32,
//! .m64n72k32, .m64n80k32, .m64n88k32, .m64n96k32,
//! .m64n104k32, .m64n112k32, .m64n120k32, .m64n128k32,
//! .m64n136k32, .m64n144k32, .m64n152k32, .m64n160k32,
//! .m64n168k32, .m64n176k32, .m64n184k32, .m64n192k32,
//! .m64n200k32, .m64n208k32, .m64n216k32, .m64n224k32,
//! .m64n232k32, .m64n240k32, .m64n248k32, .m64n256k32};
//! .dtype  = {.f32};
//! ------------------------------------------------------------------
//! // .tf32 floating point type:
//! wgmma.mma_async.sp.sync.aligned.shape.dtype.tf32.tf32  d, a-desc, b-desc, sp-meta, sp-sel, scale-d, imm-scale-a, imm-scale-b;
//! wgmma.mma_async.sp.sync.aligned.shape.dtype.tf32.tf32  d, a, b-desc, sp-meta, sp-sel, scale-d, imm-scale-a, imm-scale-b;
//! .shape   = {.m64n8k16, .m64n16k16, .m64n24k16, .m64n32k16,
//! .m64n40k16, .m64n48k16, .m64n56k16, .m64n64k16,
//! .m64n72k16, .m64n80k16, .m64n88k16, .m64n96k16,
//! .m64n104k16, .m64n112k16, .m64n120k16, .m64n128k16,
//! .m64n136k16, .m64n144k16, .m64n152k16, .m64n160k16,
//! .m64n168k16, .m64n176k16, .m64n184k16, .m64n192k16,
//! .m64n200k16, .m64n208k16, .m64n216k16, .m64n224k16,
//! .m64n232k16, .m64n240k16, .m64n248k16, .m64n256k16};
//! .dtype  = {.f32};
//! ------------------------------------------------------------------
//! // FP8 floating point type
//! wgmma.mma_async.sp.sync.aligned.shape.dtype.atype.btype  d, a-desc, b-desc, sp-meta, sp-sel, scale-d, imm-scale-a, imm-scale-b;
//! wgmma.mma_async.sp.sync.aligned.shape.dtype.atype.btype  d, a, b-desc, sp-meta, sp-sel, scale-d, imm-scale-a, imm-scale-b;
//! .shape   = {.m64n8k64, .m64n16k64, .m64n24k64, .m64n32k64,
//! .m64n40k64, .m64n48k64, .m64n56k64, .m64n64k64,
//! .m64n72k64, .m64n80k64, .m64n88k64, .m64n96k64,
//! .m64n104k64, .m64n112k64, .m64n120k64, .m64n128k64,
//! .m64n136k64, .m64n144k64, .m64n152k64, .m64n160k64,
//! .m64n168k64, .m64n176k64, .m64n184k64, .m64n192k64,
//! .m64n200k64, .m64n208k64, .m64n216k64, .m64n224k64,
//! .m64n232k64, .m64n240k64, .m64n248k64, .m64n256k64};
//! .atype  = {.e4m3, .e5m2};
//! .btype  = {.e4m3, .e5m2};
//! .dtype  = {.f16, .f32};
//! ------------------------------------------------------------------
//! // Integer type:
//! wgmma.mma_async.sp.sync.aligned.shape{.satfinite}.s32.atype.btype  d, a-desc, b-desc, sp-meta, sp-sel, scale-d;
//! wgmma.mma_async.sp.sync.aligned.shape{.satfinite}.s32.atype.btype  d, a, b-desc, sp-meta, sp-sel, scale-d;
//! .shape   = {.m64n8k64, .m64n16k64, .m64n24k64, .m64n32k64,
//! .m64n48k64, .m64n64k64, .m64n80k64, .m64n96k64,
//! .m64n112k64, .m64n128k64, .m64n144k64, .m64n160k64,
//! .m64n176k64, .m64n192k64, .m64n208k64, .m64n224k64,
//! .m64n240k64, .m64n256k64};
//! .atype  = {.s8, .u8};
//! .btype  = {.s8, .u8};

#![allow(unused)]

use crate::lexer::PtxToken;
use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
use crate::r#type::common::*;

pub mod section_0 {
    use super::*;
    use crate::r#type::instruction::wgmma_mma_async_sp::section_0::*;

    // ============================================================================
    // Generated enum parsers
    // ============================================================================

    impl PtxParser for Dtype {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            // Try F16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".f16").is_ok() {
                    return Ok(Dtype::F16);
                }
                stream.set_position(saved_pos);
            }
            let saved_pos = stream.position();
            // Try F32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".f32").is_ok() {
                    return Ok(Dtype::F32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let span = stream
                .peek()
                .map(|(_, s)| s.clone())
                .unwrap_or(Span { start: 0, end: 0 });
            let expected = &[".f16", ".f32"];
            let found = stream
                .peek()
                .map(|(t, _)| format!("{:?}", t))
                .unwrap_or_else(|_| "<end of input>".to_string());
            Err(crate::parser::unexpected_value(span, expected, found))
        }
    }

    impl PtxParser for Shape {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            // Try M64n104k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n104k32").is_ok() {
                    return Ok(Shape::M64n104k32);
                }
                stream.set_position(saved_pos);
            }
            let saved_pos = stream.position();
            // Try M64n112k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n112k32").is_ok() {
                    return Ok(Shape::M64n112k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n120k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n120k32").is_ok() {
                    return Ok(Shape::M64n120k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n128k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n128k32").is_ok() {
                    return Ok(Shape::M64n128k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n136k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n136k32").is_ok() {
                    return Ok(Shape::M64n136k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n144k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n144k32").is_ok() {
                    return Ok(Shape::M64n144k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n152k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n152k32").is_ok() {
                    return Ok(Shape::M64n152k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n160k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n160k32").is_ok() {
                    return Ok(Shape::M64n160k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n168k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n168k32").is_ok() {
                    return Ok(Shape::M64n168k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n176k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n176k32").is_ok() {
                    return Ok(Shape::M64n176k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n184k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n184k32").is_ok() {
                    return Ok(Shape::M64n184k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n192k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n192k32").is_ok() {
                    return Ok(Shape::M64n192k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n200k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n200k32").is_ok() {
                    return Ok(Shape::M64n200k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n208k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n208k32").is_ok() {
                    return Ok(Shape::M64n208k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n216k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n216k32").is_ok() {
                    return Ok(Shape::M64n216k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n224k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n224k32").is_ok() {
                    return Ok(Shape::M64n224k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n232k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n232k32").is_ok() {
                    return Ok(Shape::M64n232k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n240k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n240k32").is_ok() {
                    return Ok(Shape::M64n240k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n248k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n248k32").is_ok() {
                    return Ok(Shape::M64n248k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n256k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n256k32").is_ok() {
                    return Ok(Shape::M64n256k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n16k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n16k32").is_ok() {
                    return Ok(Shape::M64n16k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n24k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n24k32").is_ok() {
                    return Ok(Shape::M64n24k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n32k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n32k32").is_ok() {
                    return Ok(Shape::M64n32k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n40k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n40k32").is_ok() {
                    return Ok(Shape::M64n40k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n48k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n48k32").is_ok() {
                    return Ok(Shape::M64n48k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n56k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n56k32").is_ok() {
                    return Ok(Shape::M64n56k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n64k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n64k32").is_ok() {
                    return Ok(Shape::M64n64k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n72k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n72k32").is_ok() {
                    return Ok(Shape::M64n72k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n80k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n80k32").is_ok() {
                    return Ok(Shape::M64n80k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n88k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n88k32").is_ok() {
                    return Ok(Shape::M64n88k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n96k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n96k32").is_ok() {
                    return Ok(Shape::M64n96k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n8k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n8k32").is_ok() {
                    return Ok(Shape::M64n8k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let span = stream
                .peek()
                .map(|(_, s)| s.clone())
                .unwrap_or(Span { start: 0, end: 0 });
            let expected = &[
                ".m64n104k32",
                ".m64n112k32",
                ".m64n120k32",
                ".m64n128k32",
                ".m64n136k32",
                ".m64n144k32",
                ".m64n152k32",
                ".m64n160k32",
                ".m64n168k32",
                ".m64n176k32",
                ".m64n184k32",
                ".m64n192k32",
                ".m64n200k32",
                ".m64n208k32",
                ".m64n216k32",
                ".m64n224k32",
                ".m64n232k32",
                ".m64n240k32",
                ".m64n248k32",
                ".m64n256k32",
                ".m64n16k32",
                ".m64n24k32",
                ".m64n32k32",
                ".m64n40k32",
                ".m64n48k32",
                ".m64n56k32",
                ".m64n64k32",
                ".m64n72k32",
                ".m64n80k32",
                ".m64n88k32",
                ".m64n96k32",
                ".m64n8k32",
            ];
            let found = stream
                .peek()
                .map(|(t, _)| format!("{:?}", t))
                .unwrap_or_else(|_| "<end of input>".to_string());
            Err(crate::parser::unexpected_value(span, expected, found))
        }
    }

    impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeDtypeF16F16 {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            stream.expect_string("wgmma")?;
            stream.expect_string(".mma_async")?;
            let mma_async = ();
            stream.expect_complete()?;
            stream.expect_string(".sp")?;
            let sp = ();
            stream.expect_complete()?;
            stream.expect_string(".sync")?;
            let sync = ();
            stream.expect_complete()?;
            stream.expect_string(".aligned")?;
            let aligned = ();
            stream.expect_complete()?;
            let shape = Shape::parse(stream)?;
            stream.expect_complete()?;
            let dtype = Dtype::parse(stream)?;
            stream.expect_complete()?;
            stream.expect_string(".f16")?;
            let f16 = ();
            stream.expect_complete()?;
            stream.expect_string(".f16")?;
            let f162 = ();
            stream.expect_complete()?;
            let d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let a_desc = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let b_desc = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_meta = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_sel = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let scale_d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_scale_a = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_scale_b = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_trans_a = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_trans_b = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Semicolon)?;
            Ok(WgmmaMmaAsyncSpSyncAlignedShapeDtypeF16F16 {
                mma_async,
                sp,
                sync,
                aligned,
                shape,
                dtype,
                f16,
                f162,
                d,
                a_desc,
                b_desc,
                sp_meta,
                sp_sel,
                scale_d,
                imm_scale_a,
                imm_scale_b,
                imm_trans_a,
                imm_trans_b,
            })
        }
    }

    impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeDtypeF16F161 {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            stream.expect_string("wgmma")?;
            stream.expect_string(".mma_async")?;
            let mma_async = ();
            stream.expect_complete()?;
            stream.expect_string(".sp")?;
            let sp = ();
            stream.expect_complete()?;
            stream.expect_string(".sync")?;
            let sync = ();
            stream.expect_complete()?;
            stream.expect_string(".aligned")?;
            let aligned = ();
            stream.expect_complete()?;
            let shape = Shape::parse(stream)?;
            stream.expect_complete()?;
            let dtype = Dtype::parse(stream)?;
            stream.expect_complete()?;
            stream.expect_string(".f16")?;
            let f16 = ();
            stream.expect_complete()?;
            stream.expect_string(".f16")?;
            let f162 = ();
            stream.expect_complete()?;
            let d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let a = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let b_desc = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_meta = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_sel = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let scale_d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_scale_a = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_scale_b = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_trans_b = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Semicolon)?;
            Ok(WgmmaMmaAsyncSpSyncAlignedShapeDtypeF16F161 {
                mma_async,
                sp,
                sync,
                aligned,
                shape,
                dtype,
                f16,
                f162,
                d,
                a,
                b_desc,
                sp_meta,
                sp_sel,
                scale_d,
                imm_scale_a,
                imm_scale_b,
                imm_trans_b,
            })
        }
    }
}

pub mod section_1 {
    use super::*;
    use crate::r#type::instruction::wgmma_mma_async_sp::section_1::*;

    // ============================================================================
    // Generated enum parsers
    // ============================================================================

    impl PtxParser for Dtype {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            // Try F32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".f32").is_ok() {
                    return Ok(Dtype::F32);
                }
                stream.set_position(saved_pos);
            }
            let span = stream
                .peek()
                .map(|(_, s)| s.clone())
                .unwrap_or(Span { start: 0, end: 0 });
            let expected = &[".f32"];
            let found = stream
                .peek()
                .map(|(t, _)| format!("{:?}", t))
                .unwrap_or_else(|_| "<end of input>".to_string());
            Err(crate::parser::unexpected_value(span, expected, found))
        }
    }

    impl PtxParser for Shape {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            // Try M64n104k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n104k32").is_ok() {
                    return Ok(Shape::M64n104k32);
                }
                stream.set_position(saved_pos);
            }
            let saved_pos = stream.position();
            // Try M64n112k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n112k32").is_ok() {
                    return Ok(Shape::M64n112k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n120k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n120k32").is_ok() {
                    return Ok(Shape::M64n120k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n128k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n128k32").is_ok() {
                    return Ok(Shape::M64n128k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n136k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n136k32").is_ok() {
                    return Ok(Shape::M64n136k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n144k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n144k32").is_ok() {
                    return Ok(Shape::M64n144k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n152k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n152k32").is_ok() {
                    return Ok(Shape::M64n152k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n160k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n160k32").is_ok() {
                    return Ok(Shape::M64n160k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n168k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n168k32").is_ok() {
                    return Ok(Shape::M64n168k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n176k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n176k32").is_ok() {
                    return Ok(Shape::M64n176k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n184k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n184k32").is_ok() {
                    return Ok(Shape::M64n184k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n192k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n192k32").is_ok() {
                    return Ok(Shape::M64n192k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n200k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n200k32").is_ok() {
                    return Ok(Shape::M64n200k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n208k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n208k32").is_ok() {
                    return Ok(Shape::M64n208k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n216k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n216k32").is_ok() {
                    return Ok(Shape::M64n216k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n224k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n224k32").is_ok() {
                    return Ok(Shape::M64n224k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n232k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n232k32").is_ok() {
                    return Ok(Shape::M64n232k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n240k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n240k32").is_ok() {
                    return Ok(Shape::M64n240k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n248k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n248k32").is_ok() {
                    return Ok(Shape::M64n248k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n256k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n256k32").is_ok() {
                    return Ok(Shape::M64n256k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n16k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n16k32").is_ok() {
                    return Ok(Shape::M64n16k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n24k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n24k32").is_ok() {
                    return Ok(Shape::M64n24k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n32k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n32k32").is_ok() {
                    return Ok(Shape::M64n32k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n40k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n40k32").is_ok() {
                    return Ok(Shape::M64n40k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n48k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n48k32").is_ok() {
                    return Ok(Shape::M64n48k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n56k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n56k32").is_ok() {
                    return Ok(Shape::M64n56k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n64k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n64k32").is_ok() {
                    return Ok(Shape::M64n64k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n72k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n72k32").is_ok() {
                    return Ok(Shape::M64n72k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n80k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n80k32").is_ok() {
                    return Ok(Shape::M64n80k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n88k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n88k32").is_ok() {
                    return Ok(Shape::M64n88k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n96k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n96k32").is_ok() {
                    return Ok(Shape::M64n96k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n8k32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n8k32").is_ok() {
                    return Ok(Shape::M64n8k32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let span = stream
                .peek()
                .map(|(_, s)| s.clone())
                .unwrap_or(Span { start: 0, end: 0 });
            let expected = &[
                ".m64n104k32",
                ".m64n112k32",
                ".m64n120k32",
                ".m64n128k32",
                ".m64n136k32",
                ".m64n144k32",
                ".m64n152k32",
                ".m64n160k32",
                ".m64n168k32",
                ".m64n176k32",
                ".m64n184k32",
                ".m64n192k32",
                ".m64n200k32",
                ".m64n208k32",
                ".m64n216k32",
                ".m64n224k32",
                ".m64n232k32",
                ".m64n240k32",
                ".m64n248k32",
                ".m64n256k32",
                ".m64n16k32",
                ".m64n24k32",
                ".m64n32k32",
                ".m64n40k32",
                ".m64n48k32",
                ".m64n56k32",
                ".m64n64k32",
                ".m64n72k32",
                ".m64n80k32",
                ".m64n88k32",
                ".m64n96k32",
                ".m64n8k32",
            ];
            let found = stream
                .peek()
                .map(|(t, _)| format!("{:?}", t))
                .unwrap_or_else(|_| "<end of input>".to_string());
            Err(crate::parser::unexpected_value(span, expected, found))
        }
    }

    impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeDtypeBf16Bf16 {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            stream.expect_string("wgmma")?;
            stream.expect_string(".mma_async")?;
            let mma_async = ();
            stream.expect_complete()?;
            stream.expect_string(".sp")?;
            let sp = ();
            stream.expect_complete()?;
            stream.expect_string(".sync")?;
            let sync = ();
            stream.expect_complete()?;
            stream.expect_string(".aligned")?;
            let aligned = ();
            stream.expect_complete()?;
            let shape = Shape::parse(stream)?;
            stream.expect_complete()?;
            let dtype = Dtype::parse(stream)?;
            stream.expect_complete()?;
            stream.expect_string(".bf16")?;
            let bf16 = ();
            stream.expect_complete()?;
            stream.expect_string(".bf16")?;
            let bf162 = ();
            stream.expect_complete()?;
            let d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let a_desc = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let b_desc = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_meta = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_sel = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let scale_d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_scale_a = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_scale_b = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_trans_a = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_trans_b = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Semicolon)?;
            Ok(WgmmaMmaAsyncSpSyncAlignedShapeDtypeBf16Bf16 {
                mma_async,
                sp,
                sync,
                aligned,
                shape,
                dtype,
                bf16,
                bf162,
                d,
                a_desc,
                b_desc,
                sp_meta,
                sp_sel,
                scale_d,
                imm_scale_a,
                imm_scale_b,
                imm_trans_a,
                imm_trans_b,
            })
        }
    }

    impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeDtypeBf16Bf161 {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            stream.expect_string("wgmma")?;
            stream.expect_string(".mma_async")?;
            let mma_async = ();
            stream.expect_complete()?;
            stream.expect_string(".sp")?;
            let sp = ();
            stream.expect_complete()?;
            stream.expect_string(".sync")?;
            let sync = ();
            stream.expect_complete()?;
            stream.expect_string(".aligned")?;
            let aligned = ();
            stream.expect_complete()?;
            let shape = Shape::parse(stream)?;
            stream.expect_complete()?;
            let dtype = Dtype::parse(stream)?;
            stream.expect_complete()?;
            stream.expect_string(".bf16")?;
            let bf16 = ();
            stream.expect_complete()?;
            stream.expect_string(".bf16")?;
            let bf162 = ();
            stream.expect_complete()?;
            let d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let a = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let b_desc = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_meta = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_sel = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let scale_d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_scale_a = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_scale_b = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_trans_b = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Semicolon)?;
            Ok(WgmmaMmaAsyncSpSyncAlignedShapeDtypeBf16Bf161 {
                mma_async,
                sp,
                sync,
                aligned,
                shape,
                dtype,
                bf16,
                bf162,
                d,
                a,
                b_desc,
                sp_meta,
                sp_sel,
                scale_d,
                imm_scale_a,
                imm_scale_b,
                imm_trans_b,
            })
        }
    }
}

pub mod section_2 {
    use super::*;
    use crate::r#type::instruction::wgmma_mma_async_sp::section_2::*;

    // ============================================================================
    // Generated enum parsers
    // ============================================================================

    impl PtxParser for Dtype {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            // Try F32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".f32").is_ok() {
                    return Ok(Dtype::F32);
                }
                stream.set_position(saved_pos);
            }
            let span = stream
                .peek()
                .map(|(_, s)| s.clone())
                .unwrap_or(Span { start: 0, end: 0 });
            let expected = &[".f32"];
            let found = stream
                .peek()
                .map(|(t, _)| format!("{:?}", t))
                .unwrap_or_else(|_| "<end of input>".to_string());
            Err(crate::parser::unexpected_value(span, expected, found))
        }
    }

    impl PtxParser for Shape {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            // Try M64n104k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n104k16").is_ok() {
                    return Ok(Shape::M64n104k16);
                }
                stream.set_position(saved_pos);
            }
            let saved_pos = stream.position();
            // Try M64n112k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n112k16").is_ok() {
                    return Ok(Shape::M64n112k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n120k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n120k16").is_ok() {
                    return Ok(Shape::M64n120k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n128k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n128k16").is_ok() {
                    return Ok(Shape::M64n128k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n136k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n136k16").is_ok() {
                    return Ok(Shape::M64n136k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n144k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n144k16").is_ok() {
                    return Ok(Shape::M64n144k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n152k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n152k16").is_ok() {
                    return Ok(Shape::M64n152k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n160k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n160k16").is_ok() {
                    return Ok(Shape::M64n160k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n168k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n168k16").is_ok() {
                    return Ok(Shape::M64n168k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n176k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n176k16").is_ok() {
                    return Ok(Shape::M64n176k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n184k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n184k16").is_ok() {
                    return Ok(Shape::M64n184k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n192k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n192k16").is_ok() {
                    return Ok(Shape::M64n192k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n200k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n200k16").is_ok() {
                    return Ok(Shape::M64n200k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n208k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n208k16").is_ok() {
                    return Ok(Shape::M64n208k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n216k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n216k16").is_ok() {
                    return Ok(Shape::M64n216k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n224k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n224k16").is_ok() {
                    return Ok(Shape::M64n224k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n232k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n232k16").is_ok() {
                    return Ok(Shape::M64n232k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n240k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n240k16").is_ok() {
                    return Ok(Shape::M64n240k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n248k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n248k16").is_ok() {
                    return Ok(Shape::M64n248k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n256k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n256k16").is_ok() {
                    return Ok(Shape::M64n256k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n16k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n16k16").is_ok() {
                    return Ok(Shape::M64n16k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n24k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n24k16").is_ok() {
                    return Ok(Shape::M64n24k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n32k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n32k16").is_ok() {
                    return Ok(Shape::M64n32k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n40k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n40k16").is_ok() {
                    return Ok(Shape::M64n40k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n48k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n48k16").is_ok() {
                    return Ok(Shape::M64n48k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n56k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n56k16").is_ok() {
                    return Ok(Shape::M64n56k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n64k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n64k16").is_ok() {
                    return Ok(Shape::M64n64k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n72k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n72k16").is_ok() {
                    return Ok(Shape::M64n72k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n80k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n80k16").is_ok() {
                    return Ok(Shape::M64n80k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n88k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n88k16").is_ok() {
                    return Ok(Shape::M64n88k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n96k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n96k16").is_ok() {
                    return Ok(Shape::M64n96k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n8k16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n8k16").is_ok() {
                    return Ok(Shape::M64n8k16);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let span = stream
                .peek()
                .map(|(_, s)| s.clone())
                .unwrap_or(Span { start: 0, end: 0 });
            let expected = &[
                ".m64n104k16",
                ".m64n112k16",
                ".m64n120k16",
                ".m64n128k16",
                ".m64n136k16",
                ".m64n144k16",
                ".m64n152k16",
                ".m64n160k16",
                ".m64n168k16",
                ".m64n176k16",
                ".m64n184k16",
                ".m64n192k16",
                ".m64n200k16",
                ".m64n208k16",
                ".m64n216k16",
                ".m64n224k16",
                ".m64n232k16",
                ".m64n240k16",
                ".m64n248k16",
                ".m64n256k16",
                ".m64n16k16",
                ".m64n24k16",
                ".m64n32k16",
                ".m64n40k16",
                ".m64n48k16",
                ".m64n56k16",
                ".m64n64k16",
                ".m64n72k16",
                ".m64n80k16",
                ".m64n88k16",
                ".m64n96k16",
                ".m64n8k16",
            ];
            let found = stream
                .peek()
                .map(|(t, _)| format!("{:?}", t))
                .unwrap_or_else(|_| "<end of input>".to_string());
            Err(crate::parser::unexpected_value(span, expected, found))
        }
    }

    impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeDtypeTf32Tf32 {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            stream.expect_string("wgmma")?;
            stream.expect_string(".mma_async")?;
            let mma_async = ();
            stream.expect_complete()?;
            stream.expect_string(".sp")?;
            let sp = ();
            stream.expect_complete()?;
            stream.expect_string(".sync")?;
            let sync = ();
            stream.expect_complete()?;
            stream.expect_string(".aligned")?;
            let aligned = ();
            stream.expect_complete()?;
            let shape = Shape::parse(stream)?;
            stream.expect_complete()?;
            let dtype = Dtype::parse(stream)?;
            stream.expect_complete()?;
            stream.expect_string(".tf32")?;
            let tf32 = ();
            stream.expect_complete()?;
            stream.expect_string(".tf32")?;
            let tf322 = ();
            stream.expect_complete()?;
            let d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let a_desc = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let b_desc = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_meta = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_sel = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let scale_d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_scale_a = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_scale_b = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Semicolon)?;
            Ok(WgmmaMmaAsyncSpSyncAlignedShapeDtypeTf32Tf32 {
                mma_async,
                sp,
                sync,
                aligned,
                shape,
                dtype,
                tf32,
                tf322,
                d,
                a_desc,
                b_desc,
                sp_meta,
                sp_sel,
                scale_d,
                imm_scale_a,
                imm_scale_b,
            })
        }
    }

    impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeDtypeTf32Tf321 {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            stream.expect_string("wgmma")?;
            stream.expect_string(".mma_async")?;
            let mma_async = ();
            stream.expect_complete()?;
            stream.expect_string(".sp")?;
            let sp = ();
            stream.expect_complete()?;
            stream.expect_string(".sync")?;
            let sync = ();
            stream.expect_complete()?;
            stream.expect_string(".aligned")?;
            let aligned = ();
            stream.expect_complete()?;
            let shape = Shape::parse(stream)?;
            stream.expect_complete()?;
            let dtype = Dtype::parse(stream)?;
            stream.expect_complete()?;
            stream.expect_string(".tf32")?;
            let tf32 = ();
            stream.expect_complete()?;
            stream.expect_string(".tf32")?;
            let tf322 = ();
            stream.expect_complete()?;
            let d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let a = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let b_desc = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_meta = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_sel = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let scale_d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_scale_a = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_scale_b = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Semicolon)?;
            Ok(WgmmaMmaAsyncSpSyncAlignedShapeDtypeTf32Tf321 {
                mma_async,
                sp,
                sync,
                aligned,
                shape,
                dtype,
                tf32,
                tf322,
                d,
                a,
                b_desc,
                sp_meta,
                sp_sel,
                scale_d,
                imm_scale_a,
                imm_scale_b,
            })
        }
    }
}

pub mod section_3 {
    use super::*;
    use crate::r#type::instruction::wgmma_mma_async_sp::section_3::*;

    // ============================================================================
    // Generated enum parsers
    // ============================================================================

    impl PtxParser for Atype {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            // Try E4m3
            {
                let saved_pos = stream.position();
                if stream.expect_string(".e4m3").is_ok() {
                    return Ok(Atype::E4m3);
                }
                stream.set_position(saved_pos);
            }
            let saved_pos = stream.position();
            // Try E5m2
            {
                let saved_pos = stream.position();
                if stream.expect_string(".e5m2").is_ok() {
                    return Ok(Atype::E5m2);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let span = stream
                .peek()
                .map(|(_, s)| s.clone())
                .unwrap_or(Span { start: 0, end: 0 });
            let expected = &[".e4m3", ".e5m2"];
            let found = stream
                .peek()
                .map(|(t, _)| format!("{:?}", t))
                .unwrap_or_else(|_| "<end of input>".to_string());
            Err(crate::parser::unexpected_value(span, expected, found))
        }
    }

    impl PtxParser for Btype {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            // Try E4m3
            {
                let saved_pos = stream.position();
                if stream.expect_string(".e4m3").is_ok() {
                    return Ok(Btype::E4m3);
                }
                stream.set_position(saved_pos);
            }
            let saved_pos = stream.position();
            // Try E5m2
            {
                let saved_pos = stream.position();
                if stream.expect_string(".e5m2").is_ok() {
                    return Ok(Btype::E5m2);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let span = stream
                .peek()
                .map(|(_, s)| s.clone())
                .unwrap_or(Span { start: 0, end: 0 });
            let expected = &[".e4m3", ".e5m2"];
            let found = stream
                .peek()
                .map(|(t, _)| format!("{:?}", t))
                .unwrap_or_else(|_| "<end of input>".to_string());
            Err(crate::parser::unexpected_value(span, expected, found))
        }
    }

    impl PtxParser for Dtype {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            // Try F16
            {
                let saved_pos = stream.position();
                if stream.expect_string(".f16").is_ok() {
                    return Ok(Dtype::F16);
                }
                stream.set_position(saved_pos);
            }
            let saved_pos = stream.position();
            // Try F32
            {
                let saved_pos = stream.position();
                if stream.expect_string(".f32").is_ok() {
                    return Ok(Dtype::F32);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let span = stream
                .peek()
                .map(|(_, s)| s.clone())
                .unwrap_or(Span { start: 0, end: 0 });
            let expected = &[".f16", ".f32"];
            let found = stream
                .peek()
                .map(|(t, _)| format!("{:?}", t))
                .unwrap_or_else(|_| "<end of input>".to_string());
            Err(crate::parser::unexpected_value(span, expected, found))
        }
    }

    impl PtxParser for Shape {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            // Try M64n104k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n104k64").is_ok() {
                    return Ok(Shape::M64n104k64);
                }
                stream.set_position(saved_pos);
            }
            let saved_pos = stream.position();
            // Try M64n112k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n112k64").is_ok() {
                    return Ok(Shape::M64n112k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n120k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n120k64").is_ok() {
                    return Ok(Shape::M64n120k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n128k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n128k64").is_ok() {
                    return Ok(Shape::M64n128k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n136k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n136k64").is_ok() {
                    return Ok(Shape::M64n136k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n144k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n144k64").is_ok() {
                    return Ok(Shape::M64n144k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n152k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n152k64").is_ok() {
                    return Ok(Shape::M64n152k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n160k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n160k64").is_ok() {
                    return Ok(Shape::M64n160k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n168k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n168k64").is_ok() {
                    return Ok(Shape::M64n168k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n176k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n176k64").is_ok() {
                    return Ok(Shape::M64n176k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n184k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n184k64").is_ok() {
                    return Ok(Shape::M64n184k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n192k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n192k64").is_ok() {
                    return Ok(Shape::M64n192k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n200k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n200k64").is_ok() {
                    return Ok(Shape::M64n200k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n208k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n208k64").is_ok() {
                    return Ok(Shape::M64n208k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n216k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n216k64").is_ok() {
                    return Ok(Shape::M64n216k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n224k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n224k64").is_ok() {
                    return Ok(Shape::M64n224k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n232k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n232k64").is_ok() {
                    return Ok(Shape::M64n232k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n240k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n240k64").is_ok() {
                    return Ok(Shape::M64n240k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n248k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n248k64").is_ok() {
                    return Ok(Shape::M64n248k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n256k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n256k64").is_ok() {
                    return Ok(Shape::M64n256k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n16k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n16k64").is_ok() {
                    return Ok(Shape::M64n16k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n24k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n24k64").is_ok() {
                    return Ok(Shape::M64n24k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n32k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n32k64").is_ok() {
                    return Ok(Shape::M64n32k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n40k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n40k64").is_ok() {
                    return Ok(Shape::M64n40k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n48k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n48k64").is_ok() {
                    return Ok(Shape::M64n48k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n56k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n56k64").is_ok() {
                    return Ok(Shape::M64n56k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n64k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n64k64").is_ok() {
                    return Ok(Shape::M64n64k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n72k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n72k64").is_ok() {
                    return Ok(Shape::M64n72k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n80k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n80k64").is_ok() {
                    return Ok(Shape::M64n80k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n88k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n88k64").is_ok() {
                    return Ok(Shape::M64n88k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n96k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n96k64").is_ok() {
                    return Ok(Shape::M64n96k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n8k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n8k64").is_ok() {
                    return Ok(Shape::M64n8k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let span = stream
                .peek()
                .map(|(_, s)| s.clone())
                .unwrap_or(Span { start: 0, end: 0 });
            let expected = &[
                ".m64n104k64",
                ".m64n112k64",
                ".m64n120k64",
                ".m64n128k64",
                ".m64n136k64",
                ".m64n144k64",
                ".m64n152k64",
                ".m64n160k64",
                ".m64n168k64",
                ".m64n176k64",
                ".m64n184k64",
                ".m64n192k64",
                ".m64n200k64",
                ".m64n208k64",
                ".m64n216k64",
                ".m64n224k64",
                ".m64n232k64",
                ".m64n240k64",
                ".m64n248k64",
                ".m64n256k64",
                ".m64n16k64",
                ".m64n24k64",
                ".m64n32k64",
                ".m64n40k64",
                ".m64n48k64",
                ".m64n56k64",
                ".m64n64k64",
                ".m64n72k64",
                ".m64n80k64",
                ".m64n88k64",
                ".m64n96k64",
                ".m64n8k64",
            ];
            let found = stream
                .peek()
                .map(|(t, _)| format!("{:?}", t))
                .unwrap_or_else(|_| "<end of input>".to_string());
            Err(crate::parser::unexpected_value(span, expected, found))
        }
    }

    impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeDtypeAtypeBtype {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            stream.expect_string("wgmma")?;
            stream.expect_string(".mma_async")?;
            let mma_async = ();
            stream.expect_complete()?;
            stream.expect_string(".sp")?;
            let sp = ();
            stream.expect_complete()?;
            stream.expect_string(".sync")?;
            let sync = ();
            stream.expect_complete()?;
            stream.expect_string(".aligned")?;
            let aligned = ();
            stream.expect_complete()?;
            let shape = Shape::parse(stream)?;
            stream.expect_complete()?;
            let dtype = Dtype::parse(stream)?;
            stream.expect_complete()?;
            let atype = Atype::parse(stream)?;
            stream.expect_complete()?;
            let btype = Btype::parse(stream)?;
            stream.expect_complete()?;
            let d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let a_desc = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let b_desc = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_meta = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_sel = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let scale_d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_scale_a = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_scale_b = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Semicolon)?;
            Ok(WgmmaMmaAsyncSpSyncAlignedShapeDtypeAtypeBtype {
                mma_async,
                sp,
                sync,
                aligned,
                shape,
                dtype,
                atype,
                btype,
                d,
                a_desc,
                b_desc,
                sp_meta,
                sp_sel,
                scale_d,
                imm_scale_a,
                imm_scale_b,
            })
        }
    }

    impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeDtypeAtypeBtype1 {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            stream.expect_string("wgmma")?;
            stream.expect_string(".mma_async")?;
            let mma_async = ();
            stream.expect_complete()?;
            stream.expect_string(".sp")?;
            let sp = ();
            stream.expect_complete()?;
            stream.expect_string(".sync")?;
            let sync = ();
            stream.expect_complete()?;
            stream.expect_string(".aligned")?;
            let aligned = ();
            stream.expect_complete()?;
            let shape = Shape::parse(stream)?;
            stream.expect_complete()?;
            let dtype = Dtype::parse(stream)?;
            stream.expect_complete()?;
            let atype = Atype::parse(stream)?;
            stream.expect_complete()?;
            let btype = Btype::parse(stream)?;
            stream.expect_complete()?;
            let d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let a = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let b_desc = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_meta = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_sel = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let scale_d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_scale_a = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let imm_scale_b = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Semicolon)?;
            Ok(WgmmaMmaAsyncSpSyncAlignedShapeDtypeAtypeBtype1 {
                mma_async,
                sp,
                sync,
                aligned,
                shape,
                dtype,
                atype,
                btype,
                d,
                a,
                b_desc,
                sp_meta,
                sp_sel,
                scale_d,
                imm_scale_a,
                imm_scale_b,
            })
        }
    }
}

pub mod section_4 {
    use super::*;
    use crate::r#type::instruction::wgmma_mma_async_sp::section_4::*;

    // ============================================================================
    // Generated enum parsers
    // ============================================================================

    impl PtxParser for Atype {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            // Try S8
            {
                let saved_pos = stream.position();
                if stream.expect_string(".s8").is_ok() {
                    return Ok(Atype::S8);
                }
                stream.set_position(saved_pos);
            }
            let saved_pos = stream.position();
            // Try U8
            {
                let saved_pos = stream.position();
                if stream.expect_string(".u8").is_ok() {
                    return Ok(Atype::U8);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let span = stream
                .peek()
                .map(|(_, s)| s.clone())
                .unwrap_or(Span { start: 0, end: 0 });
            let expected = &[".s8", ".u8"];
            let found = stream
                .peek()
                .map(|(t, _)| format!("{:?}", t))
                .unwrap_or_else(|_| "<end of input>".to_string());
            Err(crate::parser::unexpected_value(span, expected, found))
        }
    }

    impl PtxParser for Btype {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            // Try S8
            {
                let saved_pos = stream.position();
                if stream.expect_string(".s8").is_ok() {
                    return Ok(Btype::S8);
                }
                stream.set_position(saved_pos);
            }
            let saved_pos = stream.position();
            // Try U8
            {
                let saved_pos = stream.position();
                if stream.expect_string(".u8").is_ok() {
                    return Ok(Btype::U8);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let span = stream
                .peek()
                .map(|(_, s)| s.clone())
                .unwrap_or(Span { start: 0, end: 0 });
            let expected = &[".s8", ".u8"];
            let found = stream
                .peek()
                .map(|(t, _)| format!("{:?}", t))
                .unwrap_or_else(|_| "<end of input>".to_string());
            Err(crate::parser::unexpected_value(span, expected, found))
        }
    }

    impl PtxParser for Shape {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            // Try M64n112k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n112k64").is_ok() {
                    return Ok(Shape::M64n112k64);
                }
                stream.set_position(saved_pos);
            }
            let saved_pos = stream.position();
            // Try M64n128k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n128k64").is_ok() {
                    return Ok(Shape::M64n128k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n144k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n144k64").is_ok() {
                    return Ok(Shape::M64n144k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n160k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n160k64").is_ok() {
                    return Ok(Shape::M64n160k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n176k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n176k64").is_ok() {
                    return Ok(Shape::M64n176k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n192k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n192k64").is_ok() {
                    return Ok(Shape::M64n192k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n208k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n208k64").is_ok() {
                    return Ok(Shape::M64n208k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n224k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n224k64").is_ok() {
                    return Ok(Shape::M64n224k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n240k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n240k64").is_ok() {
                    return Ok(Shape::M64n240k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n256k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n256k64").is_ok() {
                    return Ok(Shape::M64n256k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n16k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n16k64").is_ok() {
                    return Ok(Shape::M64n16k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n24k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n24k64").is_ok() {
                    return Ok(Shape::M64n24k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n32k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n32k64").is_ok() {
                    return Ok(Shape::M64n32k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n48k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n48k64").is_ok() {
                    return Ok(Shape::M64n48k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n64k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n64k64").is_ok() {
                    return Ok(Shape::M64n64k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n80k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n80k64").is_ok() {
                    return Ok(Shape::M64n80k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n96k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n96k64").is_ok() {
                    return Ok(Shape::M64n96k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let saved_pos = stream.position();
            // Try M64n8k64
            {
                let saved_pos = stream.position();
                if stream.expect_string(".m64n8k64").is_ok() {
                    return Ok(Shape::M64n8k64);
                }
                stream.set_position(saved_pos);
            }
            stream.set_position(saved_pos);
            let span = stream
                .peek()
                .map(|(_, s)| s.clone())
                .unwrap_or(Span { start: 0, end: 0 });
            let expected = &[
                ".m64n112k64",
                ".m64n128k64",
                ".m64n144k64",
                ".m64n160k64",
                ".m64n176k64",
                ".m64n192k64",
                ".m64n208k64",
                ".m64n224k64",
                ".m64n240k64",
                ".m64n256k64",
                ".m64n16k64",
                ".m64n24k64",
                ".m64n32k64",
                ".m64n48k64",
                ".m64n64k64",
                ".m64n80k64",
                ".m64n96k64",
                ".m64n8k64",
            ];
            let found = stream
                .peek()
                .map(|(t, _)| format!("{:?}", t))
                .unwrap_or_else(|_| "<end of input>".to_string());
            Err(crate::parser::unexpected_value(span, expected, found))
        }
    }

    impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeSatfiniteS32AtypeBtype {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            stream.expect_string("wgmma")?;
            stream.expect_string(".mma_async")?;
            let mma_async = ();
            stream.expect_complete()?;
            stream.expect_string(".sp")?;
            let sp = ();
            stream.expect_complete()?;
            stream.expect_string(".sync")?;
            let sync = ();
            stream.expect_complete()?;
            stream.expect_string(".aligned")?;
            let aligned = ();
            stream.expect_complete()?;
            let shape = Shape::parse(stream)?;
            stream.expect_complete()?;
            let saved_pos = stream.position();
            let satfinite = stream.expect_string(".satfinite").is_ok();
            if !satfinite {
                stream.set_position(saved_pos);
            }
            stream.expect_complete()?;
            stream.expect_string(".s32")?;
            let s32 = ();
            stream.expect_complete()?;
            let atype = Atype::parse(stream)?;
            stream.expect_complete()?;
            let btype = Btype::parse(stream)?;
            stream.expect_complete()?;
            let d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let a_desc = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let b_desc = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_meta = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_sel = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let scale_d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Semicolon)?;
            Ok(WgmmaMmaAsyncSpSyncAlignedShapeSatfiniteS32AtypeBtype {
                mma_async,
                sp,
                sync,
                aligned,
                shape,
                satfinite,
                s32,
                atype,
                btype,
                d,
                a_desc,
                b_desc,
                sp_meta,
                sp_sel,
                scale_d,
            })
        }
    }

    impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeSatfiniteS32AtypeBtype1 {
        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
            stream.expect_string("wgmma")?;
            stream.expect_string(".mma_async")?;
            let mma_async = ();
            stream.expect_complete()?;
            stream.expect_string(".sp")?;
            let sp = ();
            stream.expect_complete()?;
            stream.expect_string(".sync")?;
            let sync = ();
            stream.expect_complete()?;
            stream.expect_string(".aligned")?;
            let aligned = ();
            stream.expect_complete()?;
            let shape = Shape::parse(stream)?;
            stream.expect_complete()?;
            let saved_pos = stream.position();
            let satfinite = stream.expect_string(".satfinite").is_ok();
            if !satfinite {
                stream.set_position(saved_pos);
            }
            stream.expect_complete()?;
            stream.expect_string(".s32")?;
            let s32 = ();
            stream.expect_complete()?;
            let atype = Atype::parse(stream)?;
            stream.expect_complete()?;
            let btype = Btype::parse(stream)?;
            stream.expect_complete()?;
            let d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let a = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let b_desc = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_meta = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let sp_sel = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Comma)?;
            let scale_d = GeneralOperand::parse(stream)?;
            stream.expect_complete()?;
            stream.expect_complete()?;
            stream.expect(&PtxToken::Semicolon)?;
            Ok(WgmmaMmaAsyncSpSyncAlignedShapeSatfiniteS32AtypeBtype1 {
                mma_async,
                sp,
                sync,
                aligned,
                shape,
                satfinite,
                s32,
                atype,
                btype,
                d,
                a,
                b_desc,
                sp_meta,
                sp_sel,
                scale_d,
            })
        }
    }
}