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

ptx_parser/parser/instruction/
mbarrier_test_wait.rs
