ptx_parser/parser/instruction/
cp_async_mbarrier_arrive.rs

1//! Original PTX specification:
2//!
3//! cp.async.mbarrier.arrive{.noinc}{.state}.b64 [addr];
4//! .state = { .shared, .shared::cta}
5
6#![allow(unused)]
7
8use crate::lexer::PtxToken;
9use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
10use crate::r#type::common::*;
11
12pub mod section_0 {
13    use super::*;
14    use crate::r#type::instruction::cp_async_mbarrier_arrive::section_0::*;
15
16    // ============================================================================
17    // Generated enum parsers
18    // ============================================================================
19
20    impl PtxParser for State {
21        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
22            // Try SharedCta
23            {
24                let saved_pos = stream.position();
25                if stream.expect_string(".shared::cta").is_ok() {
26                    return Ok(State::SharedCta);
27                }
28                stream.set_position(saved_pos);
29            }
30            let saved_pos = stream.position();
31            // Try Shared
32            {
33                let saved_pos = stream.position();
34                if stream.expect_string(".shared").is_ok() {
35                    return Ok(State::Shared);
36                }
37                stream.set_position(saved_pos);
38            }
39            stream.set_position(saved_pos);
40            let span = stream
41                .peek()
42                .map(|(_, s)| s.clone())
43                .unwrap_or(Span { start: 0, end: 0 });
44            let expected = &[".shared::cta", ".shared"];
45            let found = stream
46                .peek()
47                .map(|(t, _)| format!("{:?}", t))
48                .unwrap_or_else(|_| "<end of input>".to_string());
49            Err(crate::parser::unexpected_value(span, expected, found))
50        }
51    }
52
53    impl PtxParser for CpAsyncMbarrierArriveNoincStateB64 {
54        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
55            stream.expect_string("cp")?;
56            stream.expect_string(".async")?;
57            let async_ = ();
58            stream.expect_complete()?;
59            stream.expect_string(".mbarrier")?;
60            let mbarrier = ();
61            stream.expect_complete()?;
62            stream.expect_string(".arrive")?;
63            let arrive = ();
64            stream.expect_complete()?;
65            let saved_pos = stream.position();
66            let noinc = stream.expect_string(".noinc").is_ok();
67            if !noinc {
68                stream.set_position(saved_pos);
69            }
70            stream.expect_complete()?;
71            let saved_pos = stream.position();
72            let state = match State::parse(stream) {
73                Ok(val) => Some(val),
74                Err(_) => {
75                    stream.set_position(saved_pos);
76                    None
77                }
78            };
79            stream.expect_complete()?;
80            stream.expect_string(".b64")?;
81            let b64 = ();
82            stream.expect_complete()?;
83            let addr = AddressOperand::parse(stream)?;
84            stream.expect_complete()?;
85            stream.expect_complete()?;
86            stream.expect(&PtxToken::Semicolon)?;
87            Ok(CpAsyncMbarrierArriveNoincStateB64 {
88                async_,
89                mbarrier,
90                arrive,
91                noinc,
92                state,
93                b64,
94                addr,
95            })
96        }
97    }
98}
ptx_parser/parser/instruction/cp_async_mbarrier_arrive.rs

ptx_parser/parser/instruction/
cp_async_mbarrier_arrive.rs
