1#![allow(unused)]
13
14use crate::parser::{
15 PtxParseError, PtxParser, PtxTokenStream, Span,
16 util::{
17 between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
18 pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
19 },
20};
21use crate::r#type::common::*;
22use crate::{alt, ok, seq_n};
23
24pub mod section_0 {
25 use super::*;
26 use crate::r#type::instruction::mbarrier_arrive_drop::section_0::*;
27
28 impl PtxParser for Scope {
33 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
34 alt!(
35 map(string_p(".cluster"), |_, _span| Scope::Cluster),
36 map(string_p(".cta"), |_, _span| Scope::Cta)
37 )
38 }
39 }
40
41 impl PtxParser for Sem {
42 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
43 alt!(
44 map(string_p(".release"), |_, _span| Sem::Release),
45 map(string_p(".relaxed"), |_, _span| Sem::Relaxed)
46 )
47 }
48 }
49
50 impl PtxParser for State {
51 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
52 alt!(
53 map(string_p(".shared::cta"), |_, _span| State::SharedCta),
54 map(string_p(".shared"), |_, _span| State::Shared)
55 )
56 }
57 }
58
59 impl PtxParser for MbarrierArriveDropSemScopeStateB64 {
60 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
61 try_map(
62 seq_n!(
63 string_p("mbarrier"),
64 string_p(".arrive_drop"),
65 optional(Sem::parse()),
66 optional(Scope::parse()),
67 optional(State::parse()),
68 string_p(".b64"),
69 GeneralOperand::parse(),
70 comma_p(),
71 AddressOperand::parse(),
72 map(
73 optional(seq_n!(comma_p(), GeneralOperand::parse())),
74 |value, _| value.map(|(_, operand)| operand)
75 ),
76 semicolon_p()
77 ),
78 |(_, arrive_drop, sem, scope, state, b64, state2, _, addr, count, _), span| {
79 ok!(MbarrierArriveDropSemScopeStateB64 {
80 arrive_drop = arrive_drop,
81 sem = sem,
82 scope = scope,
83 state = state,
84 b64 = b64,
85 state2 = state2,
86 addr = addr,
87 count = count,
88
89 })
90 },
91 )
92 }
93 }
94
95 impl PtxParser for MbarrierArriveDropSemScopeSharedClusterB64 {
96 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
97 try_map(
98 seq_n!(
99 string_p("mbarrier"),
100 string_p(".arrive_drop"),
101 optional(Sem::parse()),
102 optional(Scope::parse()),
103 map(optional(string_p(".shared::cluster")), |value, _| value
104 .is_some()),
105 string_p(".b64"),
106 GeneralOperand::parse(),
107 comma_p(),
108 AddressOperand::parse(),
109 map(
110 optional(seq_n!(comma_p(), GeneralOperand::parse())),
111 |value, _| value.map(|(_, operand)| operand)
112 ),
113 semicolon_p()
114 ),
115 |(_, arrive_drop, sem, scope, shared_cluster, b64, operand, _, addr, count, _),
116 span| {
117 ok!(MbarrierArriveDropSemScopeSharedClusterB64 {
118 arrive_drop = arrive_drop,
119 sem = sem,
120 scope = scope,
121 shared_cluster = shared_cluster,
122 b64 = b64,
123 operand = operand,
124 addr = addr,
125 count = count,
126
127 })
128 },
129 )
130 }
131 }
132
133 impl PtxParser for MbarrierArriveDropExpectTxStateSemScopeB64 {
134 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
135 try_map(
136 seq_n!(
137 string_p("mbarrier"),
138 string_p(".arrive_drop"),
139 string_p(".expect_tx"),
140 optional(State::parse()),
141 optional(Sem::parse()),
142 optional(Scope::parse()),
143 string_p(".b64"),
144 GeneralOperand::parse(),
145 comma_p(),
146 AddressOperand::parse(),
147 comma_p(),
148 GeneralOperand::parse(),
149 semicolon_p()
150 ),
151 |(
152 _,
153 arrive_drop,
154 expect_tx,
155 state,
156 sem,
157 scope,
158 b64,
159 state2,
160 _,
161 addr,
162 _,
163 tx_count,
164 _,
165 ),
166 span| {
167 ok!(MbarrierArriveDropExpectTxStateSemScopeB64 {
168 arrive_drop = arrive_drop,
169 expect_tx = expect_tx,
170 state = state,
171 sem = sem,
172 scope = scope,
173 b64 = b64,
174 state2 = state2,
175 addr = addr,
176 tx_count = tx_count,
177
178 })
179 },
180 )
181 }
182 }
183
184 impl PtxParser for MbarrierArriveDropExpectTxSharedClusterSemScopeB64 {
185 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
186 try_map(
187 seq_n!(
188 string_p("mbarrier"),
189 string_p(".arrive_drop"),
190 string_p(".expect_tx"),
191 map(optional(string_p(".shared::cluster")), |value, _| value
192 .is_some()),
193 optional(Sem::parse()),
194 optional(Scope::parse()),
195 string_p(".b64"),
196 GeneralOperand::parse(),
197 comma_p(),
198 AddressOperand::parse(),
199 comma_p(),
200 GeneralOperand::parse(),
201 semicolon_p()
202 ),
203 |(
204 _,
205 arrive_drop,
206 expect_tx,
207 shared_cluster,
208 sem,
209 scope,
210 b64,
211 operand,
212 _,
213 addr,
214 _,
215 tx_count,
216 _,
217 ),
218 span| {
219 ok!(MbarrierArriveDropExpectTxSharedClusterSemScopeB64 {
220 arrive_drop = arrive_drop,
221 expect_tx = expect_tx,
222 shared_cluster = shared_cluster,
223 sem = sem,
224 scope = scope,
225 b64 = b64,
226 operand = operand,
227 addr = addr,
228 tx_count = tx_count,
229
230 })
231 },
232 )
233 }
234 }
235
236 impl PtxParser for MbarrierArriveDropNocompleteReleaseCtaStateB64 {
237 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
238 try_map(
239 seq_n!(
240 string_p("mbarrier"),
241 string_p(".arrive_drop"),
242 string_p(".noComplete"),
243 map(optional(string_p(".release")), |value, _| value.is_some()),
244 map(optional(string_p(".cta")), |value, _| value.is_some()),
245 optional(State::parse()),
246 string_p(".b64"),
247 GeneralOperand::parse(),
248 comma_p(),
249 AddressOperand::parse(),
250 comma_p(),
251 GeneralOperand::parse(),
252 semicolon_p()
253 ),
254 |(
255 _,
256 arrive_drop,
257 nocomplete,
258 release,
259 cta,
260 state,
261 b64,
262 state2,
263 _,
264 addr,
265 _,
266 count,
267 _,
268 ),
269 span| {
270 ok!(MbarrierArriveDropNocompleteReleaseCtaStateB64 {
271 arrive_drop = arrive_drop,
272 nocomplete = nocomplete,
273 release = release,
274 cta = cta,
275 state = state,
276 b64 = b64,
277 state2 = state2,
278 addr = addr,
279 count = count,
280
281 })
282 },
283 )
284 }
285 }
286}