ptx_parser/parser/instruction/
st_async.rs

1//! Original PTX specification:
2//!
3//! st.async{.sem}{.scope}{.ss}{.completion_mechanism}{.vec}.type [a], b, [mbar];
4//! .sem  =                 { .weak };
5//! .scope =                { .cluster };
6//! .ss   =                 { .shared::cluster };
7//! .type =                 { .b32, .b64,
8//! .u32, .u64,
9//! .s32, .s64,
10//! .f32, .f64 };
11//! .vec  =                 { .v2, .v4 };
12//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
13//! ---------------------------------------------------------
14//! st.async{.mmio}.sem.scope{.ss}.type [a], b;
15//! .sem =                  { .release };
16//! .scope =                { .gpu, .sys };
17//! .ss =                   { .global };
18//! .type =                 { .b8, .b16, .b32, .b64,
19//! .u8, .u16, .u32, .u64,
20//! .s8, .s16, .s32, .s64,
21//! .f32, .f64 };
22
23#![allow(unused)]
24
25use crate::parser::{
26    PtxParseError, PtxParser, PtxTokenStream, Span,
27    util::{
28        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
29        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
30    },
31};
32use crate::r#type::common::*;
33use crate::{alt, ok, seq_n};
34
35pub mod section_0 {
36    use super::*;
37    use crate::r#type::instruction::st_async::section_0::*;
38
39    // ============================================================================
40    // Generated enum parsers
41    // ============================================================================
42
43    impl PtxParser for CompletionMechanism {
44        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
45            alt!(map(
46                string_p(".mbarrier::complete_tx::bytes"),
47                |_, _span| CompletionMechanism::MbarrierCompleteTxBytes
48            ))
49        }
50    }
51
52    impl PtxParser for Scope {
53        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
54            alt!(map(string_p(".cluster"), |_, _span| Scope::Cluster))
55        }
56    }
57
58    impl PtxParser for Sem {
59        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
60            alt!(map(string_p(".weak"), |_, _span| Sem::Weak))
61        }
62    }
63
64    impl PtxParser for Ss {
65        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
66            alt!(map(string_p(".shared::cluster"), |_, _span| {
67                Ss::SharedCluster
68            }))
69        }
70    }
71
72    impl PtxParser for Type {
73        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
74            alt!(
75                map(string_p(".b32"), |_, _span| Type::B32),
76                map(string_p(".b64"), |_, _span| Type::B64),
77                map(string_p(".u32"), |_, _span| Type::U32),
78                map(string_p(".u64"), |_, _span| Type::U64),
79                map(string_p(".s32"), |_, _span| Type::S32),
80                map(string_p(".s64"), |_, _span| Type::S64),
81                map(string_p(".f32"), |_, _span| Type::F32),
82                map(string_p(".f64"), |_, _span| Type::F64)
83            )
84        }
85    }
86
87    impl PtxParser for Vec {
88        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
89            alt!(
90                map(string_p(".v2"), |_, _span| Vec::V2),
91                map(string_p(".v4"), |_, _span| Vec::V4)
92            )
93        }
94    }
95
96    impl PtxParser for StAsyncSemScopeSsCompletionMechanismVecType {
97        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
98            try_map(
99                seq_n!(
100                    string_p("st"),
101                    string_p(".async"),
102                    optional(Sem::parse()),
103                    optional(Scope::parse()),
104                    optional(Ss::parse()),
105                    optional(CompletionMechanism::parse()),
106                    optional(Vec::parse()),
107                    Type::parse(),
108                    AddressOperand::parse(),
109                    comma_p(),
110                    GeneralOperand::parse(),
111                    comma_p(),
112                    AddressOperand::parse(),
113                    semicolon_p()
114                ),
115                |(
116                    _,
117                    async_,
118                    sem,
119                    scope,
120                    ss,
121                    completion_mechanism,
122                    vec,
123                    type_,
124                    a,
125                    _,
126                    b,
127                    _,
128                    mbar,
129                    _,
130                ),
131                 span| {
132                    ok!(StAsyncSemScopeSsCompletionMechanismVecType {
133                        async_ = async_,
134                        sem = sem,
135                        scope = scope,
136                        ss = ss,
137                        completion_mechanism = completion_mechanism,
138                        vec = vec,
139                        type_ = type_,
140                        a = a,
141                        b = b,
142                        mbar = mbar,
143
144                    })
145                },
146            )
147        }
148    }
149}
150
151pub mod section_1 {
152    use super::*;
153    use crate::r#type::instruction::st_async::section_1::*;
154
155    // ============================================================================
156    // Generated enum parsers
157    // ============================================================================
158
159    impl PtxParser for Scope {
160        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
161            alt!(
162                map(string_p(".gpu"), |_, _span| Scope::Gpu),
163                map(string_p(".sys"), |_, _span| Scope::Sys)
164            )
165        }
166    }
167
168    impl PtxParser for Sem {
169        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
170            alt!(map(string_p(".release"), |_, _span| Sem::Release))
171        }
172    }
173
174    impl PtxParser for Ss {
175        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
176            alt!(map(string_p(".global"), |_, _span| Ss::Global))
177        }
178    }
179
180    impl PtxParser for Type {
181        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
182            alt!(
183                map(string_p(".b16"), |_, _span| Type::B16),
184                map(string_p(".b32"), |_, _span| Type::B32),
185                map(string_p(".b64"), |_, _span| Type::B64),
186                map(string_p(".u16"), |_, _span| Type::U16),
187                map(string_p(".u32"), |_, _span| Type::U32),
188                map(string_p(".u64"), |_, _span| Type::U64),
189                map(string_p(".s16"), |_, _span| Type::S16),
190                map(string_p(".s32"), |_, _span| Type::S32),
191                map(string_p(".s64"), |_, _span| Type::S64),
192                map(string_p(".f32"), |_, _span| Type::F32),
193                map(string_p(".f64"), |_, _span| Type::F64),
194                map(string_p(".b8"), |_, _span| Type::B8),
195                map(string_p(".u8"), |_, _span| Type::U8),
196                map(string_p(".s8"), |_, _span| Type::S8)
197            )
198        }
199    }
200
201    impl PtxParser for StAsyncMmioSemScopeSsType {
202        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
203            try_map(
204                seq_n!(
205                    string_p("st"),
206                    string_p(".async"),
207                    map(optional(string_p(".mmio")), |value, _| value.is_some()),
208                    Sem::parse(),
209                    Scope::parse(),
210                    optional(Ss::parse()),
211                    Type::parse(),
212                    AddressOperand::parse(),
213                    comma_p(),
214                    GeneralOperand::parse(),
215                    semicolon_p()
216                ),
217                |(_, async_, mmio, sem, scope, ss, type_, a, _, b, _), span| {
218                    ok!(StAsyncMmioSemScopeSsType {
219                        async_ = async_,
220                        mmio = mmio,
221                        sem = sem,
222                        scope = scope,
223                        ss = ss,
224                        type_ = type_,
225                        a = a,
226                        b = b,
227
228                    })
229                },
230            )
231        }
232    }
233}
ptx_parser/parser/instruction/st_async.rs

ptx_parser/parser/instruction/
st_async.rs
