ptx_parser/parser/instruction/
tcgen05_st.rs

1//! Original PTX specification:
2//!
3//! tcgen05.st.sync.aligned.shape1.num{.unpack}.b32    [taddr], r;
4//! tcgen05.st.sync.aligned.shape2.num{.unpack}.b32    [taddr], immHalfSplitoff, r;
5//! .shape1 = { .16x64b, .16x128b, .16x256b, .32x32b };
6//! .shape2 = { .16x32bx2 };
7//! .num    = { .x1, .x2, .x4, .x8, .x16, .x32, .x64, .x128 };
8//! .unpack = { .unpack::16b };
9
10#![allow(unused)]
11
12use crate::parser::{
13    PtxParseError, PtxParser, PtxTokenStream, Span,
14    util::{
15        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
16        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
17    },
18};
19use crate::r#type::common::*;
20use crate::{alt, ok, seq_n};
21
22pub mod section_0 {
23    use super::*;
24    use crate::r#type::instruction::tcgen05_st::section_0::*;
25
26    // ============================================================================
27    // Generated enum parsers
28    // ============================================================================
29
30    impl PtxParser for Num {
31        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
32            alt!(
33                map(string_p(".x128"), |_, _span| Num::X128),
34                map(string_p(".x16"), |_, _span| Num::X16),
35                map(string_p(".x32"), |_, _span| Num::X32),
36                map(string_p(".x64"), |_, _span| Num::X64),
37                map(string_p(".x1"), |_, _span| Num::X1),
38                map(string_p(".x2"), |_, _span| Num::X2),
39                map(string_p(".x4"), |_, _span| Num::X4),
40                map(string_p(".x8"), |_, _span| Num::X8)
41            )
42        }
43    }
44
45    impl PtxParser for Shape1 {
46        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
47            alt!(
48                map(string_p(".16x128b"), |_, _span| Shape1::_16x128b),
49                map(string_p(".16x256b"), |_, _span| Shape1::_16x256b),
50                map(string_p(".16x64b"), |_, _span| Shape1::_16x64b),
51                map(string_p(".32x32b"), |_, _span| Shape1::_32x32b)
52            )
53        }
54    }
55
56    impl PtxParser for Shape2 {
57        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
58            alt!(map(string_p(".16x32bx2"), |_, _span| Shape2::_16x32bx2))
59        }
60    }
61
62    impl PtxParser for Unpack {
63        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
64            alt!(map(string_p(".unpack::16b"), |_, _span| Unpack::Unpack16b))
65        }
66    }
67
68    impl PtxParser for Tcgen05StSyncAlignedShape1NumUnpackB32 {
69        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
70            try_map(
71                seq_n!(
72                    string_p("tcgen05"),
73                    string_p(".st"),
74                    string_p(".sync"),
75                    string_p(".aligned"),
76                    Shape1::parse(),
77                    Num::parse(),
78                    optional(Unpack::parse()),
79                    string_p(".b32"),
80                    AddressOperand::parse(),
81                    comma_p(),
82                    GeneralOperand::parse(),
83                    semicolon_p()
84                ),
85                |(_, st, sync, aligned, shape1, num, unpack, b32, taddr, _, r, _), span| {
86                    ok!(Tcgen05StSyncAlignedShape1NumUnpackB32 {
87                        st = st,
88                        sync = sync,
89                        aligned = aligned,
90                        shape1 = shape1,
91                        num = num,
92                        unpack = unpack,
93                        b32 = b32,
94                        taddr = taddr,
95                        r = r,
96
97                    })
98                },
99            )
100        }
101    }
102
103    impl PtxParser for Tcgen05StSyncAlignedShape2NumUnpackB32 {
104        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
105            try_map(
106                seq_n!(
107                    string_p("tcgen05"),
108                    string_p(".st"),
109                    string_p(".sync"),
110                    string_p(".aligned"),
111                    Shape2::parse(),
112                    Num::parse(),
113                    optional(Unpack::parse()),
114                    string_p(".b32"),
115                    AddressOperand::parse(),
116                    comma_p(),
117                    GeneralOperand::parse(),
118                    comma_p(),
119                    GeneralOperand::parse(),
120                    semicolon_p()
121                ),
122                |(
123                    _,
124                    st,
125                    sync,
126                    aligned,
127                    shape2,
128                    num,
129                    unpack,
130                    b32,
131                    taddr,
132                    _,
133                    immhalfsplitoff,
134                    _,
135                    r,
136                    _,
137                ),
138                 span| {
139                    ok!(Tcgen05StSyncAlignedShape2NumUnpackB32 {
140                        st = st,
141                        sync = sync,
142                        aligned = aligned,
143                        shape2 = shape2,
144                        num = num,
145                        unpack = unpack,
146                        b32 = b32,
147                        taddr = taddr,
148                        immhalfsplitoff = immhalfsplitoff,
149                        r = r,
150
151                    })
152                },
153            )
154        }
155    }
156}
ptx_parser/parser/instruction/tcgen05_st.rs

ptx_parser/parser/instruction/
tcgen05_st.rs
