ptx_parser/type/instruction/
ldmatrix.rs

1//! Original PTX specification:
2//!
3//! ldmatrix.sync.aligned.shape.num{.trans}{.ss}.type r, [p];
4//! ldmatrix.sync.aligned.m8n16.num{.ss}.dst_fmt.src_fmt        r, [p];
5//! ldmatrix.sync.aligned.m16n16.num.trans{.ss}.dst_fmt.src_fmt r, [p];
6//! .shape   = {.m8n8, .m16n16};
7//! .num     = {.x1, .x2, .x4};
8//! .ss      = {.shared, .shared::cta};
9//! .type    = {.b16, .b8};
10//! .dst_fmt = { .b8x16 };
11//! .src_fmt = { .b6x16_p32, .b4x16_p64 };
12
13#![allow(unused)]
14use crate::r#type::common::*;
15
16pub mod section_0 {
17    use crate::Spanned;
18    use crate::parser::Span;
19    use crate::r#type::common::*;
20
21    use serde::Serialize;
22
23    #[derive(Debug, Clone, PartialEq, Serialize)]
24    pub enum Shape {
25        M16n16, // .m16n16
26        M8n8,   // .m8n8
27    }
28
29    #[derive(Debug, Clone, PartialEq, Serialize)]
30    pub enum Num {
31        X1, // .x1
32        X2, // .x2
33        X4, // .x4
34    }
35
36    #[derive(Debug, Clone, PartialEq, Serialize)]
37    pub enum Ss {
38        SharedCta, // .shared::cta
39        Shared,    // .shared
40    }
41
42    #[derive(Debug, Clone, PartialEq, Serialize)]
43    pub enum Type {
44        B16, // .b16
45        B8,  // .b8
46    }
47
48    #[derive(Debug, Clone, PartialEq, Serialize)]
49    pub enum DstFmt {
50        B8x16, // .b8x16
51    }
52
53    #[derive(Debug, Clone, PartialEq, Serialize)]
54    pub enum SrcFmt {
55        B6x16P32, // .b6x16_p32
56        B4x16P64, // .b4x16_p64
57    }
58
59    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
60    pub struct LdmatrixSyncAlignedShapeNumTransSsType {
61        pub sync: (),          // .sync
62        pub aligned: (),       // .aligned
63        pub shape: Shape,      // .shape
64        pub num: Num,          // .num
65        pub trans: bool,       // {.trans}
66        pub ss: Option<Ss>,    // {.ss}
67        pub type_: Type,       // .type
68        pub r: GeneralOperand, // r
69        pub p: AddressOperand, // [p]
70        pub span: Span,
71    }
72
73    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
74    pub struct LdmatrixSyncAlignedM8n16NumSsDstFmtSrcFmt {
75        pub sync: (),          // .sync
76        pub aligned: (),       // .aligned
77        pub m8n16: (),         // .m8n16
78        pub num: Num,          // .num
79        pub ss: Option<Ss>,    // {.ss}
80        pub dst_fmt: DstFmt,   // .dst_fmt
81        pub src_fmt: SrcFmt,   // .src_fmt
82        pub r: GeneralOperand, // r
83        pub p: AddressOperand, // [p]
84        pub span: Span,
85    }
86
87    #[derive(Debug, Clone, PartialEq, Spanned, Serialize)]
88    pub struct LdmatrixSyncAlignedM16n16NumTransSsDstFmtSrcFmt {
89        pub sync: (),          // .sync
90        pub aligned: (),       // .aligned
91        pub m16n16: (),        // .m16n16
92        pub num: Num,          // .num
93        pub trans: (),         // .trans
94        pub ss: Option<Ss>,    // {.ss}
95        pub dst_fmt: DstFmt,   // .dst_fmt
96        pub src_fmt: SrcFmt,   // .src_fmt
97        pub r: GeneralOperand, // r
98        pub p: AddressOperand, // [p]
99        pub span: Span,
100    }
101}
102
103// Re-export types with section suffixes to avoid naming conflicts
104// e.g., Type0 for section_0::Type, Type1 for section_1::Type
105pub use section_0::DstFmt as DstFmt0;
106pub use section_0::LdmatrixSyncAlignedM8n16NumSsDstFmtSrcFmt;
107pub use section_0::LdmatrixSyncAlignedM16n16NumTransSsDstFmtSrcFmt;
108pub use section_0::LdmatrixSyncAlignedShapeNumTransSsType;
109pub use section_0::Num as Num0;
110pub use section_0::Shape as Shape0;
111pub use section_0::SrcFmt as SrcFmt0;
112pub use section_0::Ss as Ss0;
113pub use section_0::Type as Type0;
ptx_parser/type/instruction/ldmatrix.rs

ptx_parser/type/instruction/
ldmatrix.rs
