ptx_parser/parser/instruction/
mbarrier_arrive.rs

1//! Original PTX specification:
2//!
3//! mbarrier.arrive{.sem}{.scope}{.state}.b64           state, [addr]{, count};
4//! mbarrier.arrive{.sem}{.scope}{.shared::cluster}.b64         _, [addr] {,count};
5//! mbarrier.arrive.expect_tx{.sem}{.scope}{.state}.b64 state, [addr], txCount;
6//! mbarrier.arrive.expect_tx{.sem}{.scope}{.shared::cluster}.b64   _, [addr], txCount;
7//! mbarrier.arrive.noComplete{.release}{.cta}{.state}.b64  state, [addr], count;
8//! .sem   = { .release, .relaxed };
9//! .scope = { .cta, .cluster };
10//! .state = { .shared, .shared::cta}
11
12#![allow(unused)]
13
14use crate::lexer::PtxToken;
15use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
16use crate::r#type::common::*;
17
18pub mod section_0 {
19    use super::*;
20    use crate::r#type::instruction::mbarrier_arrive::section_0::*;
21
22    // ============================================================================
23    // Generated enum parsers
24    // ============================================================================
25
26    impl PtxParser for Scope {
27        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
28            // Try Cluster
29            {
30                let saved_pos = stream.position();
31                if stream.expect_string(".cluster").is_ok() {
32                    return Ok(Scope::Cluster);
33                }
34                stream.set_position(saved_pos);
35            }
36            let saved_pos = stream.position();
37            // Try Cta
38            {
39                let saved_pos = stream.position();
40                if stream.expect_string(".cta").is_ok() {
41                    return Ok(Scope::Cta);
42                }
43                stream.set_position(saved_pos);
44            }
45            stream.set_position(saved_pos);
46            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
47            let expected = &[".cluster", ".cta"];
48            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
49            Err(crate::parser::unexpected_value(span, expected, found))
50        }
51    }
52
53    impl PtxParser for Sem {
54        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
55            // Try Release
56            {
57                let saved_pos = stream.position();
58                if stream.expect_string(".release").is_ok() {
59                    return Ok(Sem::Release);
60                }
61                stream.set_position(saved_pos);
62            }
63            let saved_pos = stream.position();
64            // Try Relaxed
65            {
66                let saved_pos = stream.position();
67                if stream.expect_string(".relaxed").is_ok() {
68                    return Ok(Sem::Relaxed);
69                }
70                stream.set_position(saved_pos);
71            }
72            stream.set_position(saved_pos);
73            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
74            let expected = &[".release", ".relaxed"];
75            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
76            Err(crate::parser::unexpected_value(span, expected, found))
77        }
78    }
79
80    impl PtxParser for State {
81        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
82            // Try SharedCta
83            {
84                let saved_pos = stream.position();
85                if stream.expect_string(".shared::cta").is_ok() {
86                    return Ok(State::SharedCta);
87                }
88                stream.set_position(saved_pos);
89            }
90            let saved_pos = stream.position();
91            // Try Shared
92            {
93                let saved_pos = stream.position();
94                if stream.expect_string(".shared").is_ok() {
95                    return Ok(State::Shared);
96                }
97                stream.set_position(saved_pos);
98            }
99            stream.set_position(saved_pos);
100            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
101            let expected = &[".shared::cta", ".shared"];
102            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
103            Err(crate::parser::unexpected_value(span, expected, found))
104        }
105    }
106
107    impl PtxParser for MbarrierArriveSemScopeStateB64 {
108        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
109            stream.expect_string("mbarrier")?;
110            stream.expect_string(".arrive")?;
111            let arrive = ();
112            stream.expect_complete()?;
113            let saved_pos = stream.position();
114            let sem = match Sem::parse(stream) {
115                Ok(val) => Some(val),
116                Err(_) => {
117                    stream.set_position(saved_pos);
118                    None
119                }
120            };
121            stream.expect_complete()?;
122            let saved_pos = stream.position();
123            let scope = match Scope::parse(stream) {
124                Ok(val) => Some(val),
125                Err(_) => {
126                    stream.set_position(saved_pos);
127                    None
128                }
129            };
130            stream.expect_complete()?;
131            let saved_pos = stream.position();
132            let state = match State::parse(stream) {
133                Ok(val) => Some(val),
134                Err(_) => {
135                    stream.set_position(saved_pos);
136                    None
137                }
138            };
139            stream.expect_complete()?;
140            stream.expect_string(".b64")?;
141            let b64 = ();
142            stream.expect_complete()?;
143            let state2 = GeneralOperand::parse(stream)?;
144            stream.expect_complete()?;
145            stream.expect(&PtxToken::Comma)?;
146            let addr = AddressOperand::parse(stream)?;
147            stream.expect_complete()?;
148            let saved_pos = stream.position();
149            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
150            if !has_comma {
151                stream.set_position(saved_pos);
152            }
153            let saved_pos = stream.position();
154            let count = match GeneralOperand::parse(stream) {
155                Ok(val) => Some(val),
156                Err(_) => {
157                    stream.set_position(saved_pos);
158                    None
159                }
160            };
161            stream.expect_complete()?;
162            stream.expect_complete()?;
163            stream.expect(&PtxToken::Semicolon)?;
164            Ok(MbarrierArriveSemScopeStateB64 {
165                arrive,
166                sem,
167                scope,
168                state,
169                b64,
170                state2,
171                addr,
172                count,
173            })
174        }
175    }
176
177
178    impl PtxParser for MbarrierArriveSemScopeSharedClusterB64 {
179        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
180            stream.expect_string("mbarrier")?;
181            stream.expect_string(".arrive")?;
182            let arrive = ();
183            stream.expect_complete()?;
184            let saved_pos = stream.position();
185            let sem = match Sem::parse(stream) {
186                Ok(val) => Some(val),
187                Err(_) => {
188                    stream.set_position(saved_pos);
189                    None
190                }
191            };
192            stream.expect_complete()?;
193            let saved_pos = stream.position();
194            let scope = match Scope::parse(stream) {
195                Ok(val) => Some(val),
196                Err(_) => {
197                    stream.set_position(saved_pos);
198                    None
199                }
200            };
201            stream.expect_complete()?;
202            let saved_pos = stream.position();
203            let shared_cluster = stream.expect_string(".shared::cluster").is_ok();
204            if !shared_cluster {
205                stream.set_position(saved_pos);
206            }
207            stream.expect_complete()?;
208            stream.expect_string(".b64")?;
209            let b64 = ();
210            stream.expect_complete()?;
211            let operand = GeneralOperand::parse(stream)?;
212            stream.expect_complete()?;
213            stream.expect(&PtxToken::Comma)?;
214            let addr = AddressOperand::parse(stream)?;
215            stream.expect_complete()?;
216            let saved_pos = stream.position();
217            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
218            if !has_comma {
219                stream.set_position(saved_pos);
220            }
221            let saved_pos = stream.position();
222            let count = match GeneralOperand::parse(stream) {
223                Ok(val) => Some(val),
224                Err(_) => {
225                    stream.set_position(saved_pos);
226                    None
227                }
228            };
229            stream.expect_complete()?;
230            stream.expect_complete()?;
231            stream.expect(&PtxToken::Semicolon)?;
232            Ok(MbarrierArriveSemScopeSharedClusterB64 {
233                arrive,
234                sem,
235                scope,
236                shared_cluster,
237                b64,
238                operand,
239                addr,
240                count,
241            })
242        }
243    }
244
245
246    impl PtxParser for MbarrierArriveExpectTxSemScopeStateB64 {
247        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
248            stream.expect_string("mbarrier")?;
249            stream.expect_string(".arrive")?;
250            let arrive = ();
251            stream.expect_complete()?;
252            stream.expect_string(".expect_tx")?;
253            let expect_tx = ();
254            stream.expect_complete()?;
255            let saved_pos = stream.position();
256            let sem = match Sem::parse(stream) {
257                Ok(val) => Some(val),
258                Err(_) => {
259                    stream.set_position(saved_pos);
260                    None
261                }
262            };
263            stream.expect_complete()?;
264            let saved_pos = stream.position();
265            let scope = match Scope::parse(stream) {
266                Ok(val) => Some(val),
267                Err(_) => {
268                    stream.set_position(saved_pos);
269                    None
270                }
271            };
272            stream.expect_complete()?;
273            let saved_pos = stream.position();
274            let state = match State::parse(stream) {
275                Ok(val) => Some(val),
276                Err(_) => {
277                    stream.set_position(saved_pos);
278                    None
279                }
280            };
281            stream.expect_complete()?;
282            stream.expect_string(".b64")?;
283            let b64 = ();
284            stream.expect_complete()?;
285            let state2 = GeneralOperand::parse(stream)?;
286            stream.expect_complete()?;
287            stream.expect(&PtxToken::Comma)?;
288            let addr = AddressOperand::parse(stream)?;
289            stream.expect_complete()?;
290            stream.expect(&PtxToken::Comma)?;
291            let txcount = GeneralOperand::parse(stream)?;
292            stream.expect_complete()?;
293            stream.expect_complete()?;
294            stream.expect(&PtxToken::Semicolon)?;
295            Ok(MbarrierArriveExpectTxSemScopeStateB64 {
296                arrive,
297                expect_tx,
298                sem,
299                scope,
300                state,
301                b64,
302                state2,
303                addr,
304                txcount,
305            })
306        }
307    }
308
309
310    impl PtxParser for MbarrierArriveExpectTxSemScopeSharedClusterB64 {
311        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
312            stream.expect_string("mbarrier")?;
313            stream.expect_string(".arrive")?;
314            let arrive = ();
315            stream.expect_complete()?;
316            stream.expect_string(".expect_tx")?;
317            let expect_tx = ();
318            stream.expect_complete()?;
319            let saved_pos = stream.position();
320            let sem = match Sem::parse(stream) {
321                Ok(val) => Some(val),
322                Err(_) => {
323                    stream.set_position(saved_pos);
324                    None
325                }
326            };
327            stream.expect_complete()?;
328            let saved_pos = stream.position();
329            let scope = match Scope::parse(stream) {
330                Ok(val) => Some(val),
331                Err(_) => {
332                    stream.set_position(saved_pos);
333                    None
334                }
335            };
336            stream.expect_complete()?;
337            let saved_pos = stream.position();
338            let shared_cluster = stream.expect_string(".shared::cluster").is_ok();
339            if !shared_cluster {
340                stream.set_position(saved_pos);
341            }
342            stream.expect_complete()?;
343            stream.expect_string(".b64")?;
344            let b64 = ();
345            stream.expect_complete()?;
346            let operand = GeneralOperand::parse(stream)?;
347            stream.expect_complete()?;
348            stream.expect(&PtxToken::Comma)?;
349            let addr = AddressOperand::parse(stream)?;
350            stream.expect_complete()?;
351            stream.expect(&PtxToken::Comma)?;
352            let txcount = GeneralOperand::parse(stream)?;
353            stream.expect_complete()?;
354            stream.expect_complete()?;
355            stream.expect(&PtxToken::Semicolon)?;
356            Ok(MbarrierArriveExpectTxSemScopeSharedClusterB64 {
357                arrive,
358                expect_tx,
359                sem,
360                scope,
361                shared_cluster,
362                b64,
363                operand,
364                addr,
365                txcount,
366            })
367        }
368    }
369
370
371    impl PtxParser for MbarrierArriveNocompleteReleaseCtaStateB64 {
372        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
373            stream.expect_string("mbarrier")?;
374            stream.expect_string(".arrive")?;
375            let arrive = ();
376            stream.expect_complete()?;
377            stream.expect_string(".noComplete")?;
378            let nocomplete = ();
379            stream.expect_complete()?;
380            let saved_pos = stream.position();
381            let release = stream.expect_string(".release").is_ok();
382            if !release {
383                stream.set_position(saved_pos);
384            }
385            stream.expect_complete()?;
386            let saved_pos = stream.position();
387            let cta = stream.expect_string(".cta").is_ok();
388            if !cta {
389                stream.set_position(saved_pos);
390            }
391            stream.expect_complete()?;
392            let saved_pos = stream.position();
393            let state = match State::parse(stream) {
394                Ok(val) => Some(val),
395                Err(_) => {
396                    stream.set_position(saved_pos);
397                    None
398                }
399            };
400            stream.expect_complete()?;
401            stream.expect_string(".b64")?;
402            let b64 = ();
403            stream.expect_complete()?;
404            let state2 = GeneralOperand::parse(stream)?;
405            stream.expect_complete()?;
406            stream.expect(&PtxToken::Comma)?;
407            let addr = AddressOperand::parse(stream)?;
408            stream.expect_complete()?;
409            stream.expect(&PtxToken::Comma)?;
410            let count = GeneralOperand::parse(stream)?;
411            stream.expect_complete()?;
412            stream.expect_complete()?;
413            stream.expect(&PtxToken::Semicolon)?;
414            Ok(MbarrierArriveNocompleteReleaseCtaStateB64 {
415                arrive,
416                nocomplete,
417                release,
418                cta,
419                state,
420                b64,
421                state2,
422                addr,
423                count,
424            })
425        }
426    }
427
428
429}
430
ptx_parser/parser/instruction/mbarrier_arrive.rs

ptx_parser/parser/instruction/
mbarrier_arrive.rs
