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

ptx_parser/parser/instruction/
mbarrier_arrive.rs
