ptx_parser/parser/instruction/
mbarrier_test_wait.rs

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

ptx_parser/parser/instruction/
mbarrier_test_wait.rs
