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::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 MbarrierArriveSemScopeStateB64 {
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"),
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, sem, scope, state, b64, state2, _, addr, count, _), span| {
79 ok!(MbarrierArriveSemScopeStateB64 {
80 arrive = arrive,
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 MbarrierArriveSemScopeSharedClusterB64 {
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"),
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, sem, scope, shared_cluster, b64, operand, _, addr, count, _), span| {
116 ok!(MbarrierArriveSemScopeSharedClusterB64 {
117 arrive = arrive,
118 sem = sem,
119 scope = scope,
120 shared_cluster = shared_cluster,
121 b64 = b64,
122 operand = operand,
123 addr = addr,
124 count = count,
125
126 })
127 },
128 )
129 }
130 }
131
132 impl PtxParser for MbarrierArriveExpectTxSemScopeStateB64 {
133 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
134 try_map(
135 seq_n!(
136 string_p("mbarrier"),
137 string_p(".arrive"),
138 string_p(".expect_tx"),
139 optional(Sem::parse()),
140 optional(Scope::parse()),
141 optional(State::parse()),
142 string_p(".b64"),
143 GeneralOperand::parse(),
144 comma_p(),
145 AddressOperand::parse(),
146 comma_p(),
147 GeneralOperand::parse(),
148 semicolon_p()
149 ),
150 |(_, arrive, expect_tx, sem, scope, state, b64, state2, _, addr, _, txcount, _),
151 span| {
152 ok!(MbarrierArriveExpectTxSemScopeStateB64 {
153 arrive = arrive,
154 expect_tx = expect_tx,
155 sem = sem,
156 scope = scope,
157 state = state,
158 b64 = b64,
159 state2 = state2,
160 addr = addr,
161 txcount = txcount,
162
163 })
164 },
165 )
166 }
167 }
168
169 impl PtxParser for MbarrierArriveExpectTxSemScopeSharedClusterB64 {
170 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
171 try_map(
172 seq_n!(
173 string_p("mbarrier"),
174 string_p(".arrive"),
175 string_p(".expect_tx"),
176 optional(Sem::parse()),
177 optional(Scope::parse()),
178 map(optional(string_p(".shared::cluster")), |value, _| value
179 .is_some()),
180 string_p(".b64"),
181 GeneralOperand::parse(),
182 comma_p(),
183 AddressOperand::parse(),
184 comma_p(),
185 GeneralOperand::parse(),
186 semicolon_p()
187 ),
188 |(
189 _,
190 arrive,
191 expect_tx,
192 sem,
193 scope,
194 shared_cluster,
195 b64,
196 operand,
197 _,
198 addr,
199 _,
200 txcount,
201 _,
202 ),
203 span| {
204 ok!(MbarrierArriveExpectTxSemScopeSharedClusterB64 {
205 arrive = arrive,
206 expect_tx = expect_tx,
207 sem = sem,
208 scope = scope,
209 shared_cluster = shared_cluster,
210 b64 = b64,
211 operand = operand,
212 addr = addr,
213 txcount = txcount,
214
215 })
216 },
217 )
218 }
219 }
220
221 impl PtxParser for MbarrierArriveNocompleteReleaseCtaStateB64 {
222 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
223 try_map(
224 seq_n!(
225 string_p("mbarrier"),
226 string_p(".arrive"),
227 string_p(".noComplete"),
228 map(optional(string_p(".release")), |value, _| value.is_some()),
229 map(optional(string_p(".cta")), |value, _| value.is_some()),
230 optional(State::parse()),
231 string_p(".b64"),
232 GeneralOperand::parse(),
233 comma_p(),
234 AddressOperand::parse(),
235 comma_p(),
236 GeneralOperand::parse(),
237 semicolon_p()
238 ),
239 |(
240 _,
241 arrive,
242 nocomplete,
243 release,
244 cta,
245 state,
246 b64,
247 state2,
248 _,
249 addr,
250 _,
251 count,
252 _,
253 ),
254 span| {
255 ok!(MbarrierArriveNocompleteReleaseCtaStateB64 {
256 arrive = arrive,
257 nocomplete = nocomplete,
258 release = release,
259 cta = cta,
260 state = state,
261 b64 = b64,
262 state2 = state2,
263 addr = addr,
264 count = count,
265
266 })
267 },
268 )
269 }
270 }
271}