ptx_parser/parser/instruction/
mbarrier_expect_tx.rs

1//! Original PTX specification:
2//!
3//! mbarrier.expect_tx{.sem}{.scope}{.space}.b64 [addr], txCount;
4//! .sem   = { .relaxed };
5//! .scope = { .cta, .cluster };
6//! .space = { .shared, .shared::cta, .shared::cluster };
7
8#![allow(unused)]
9
10use crate::lexer::PtxToken;
11use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
12use crate::r#type::common::*;
13
14pub mod section_0 {
15    use super::*;
16    use crate::r#type::instruction::mbarrier_expect_tx::section_0::*;
17
18    // ============================================================================
19    // Generated enum parsers
20    // ============================================================================
21
22    impl PtxParser for Scope {
23        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
24            // Try Cluster
25            {
26                let saved_pos = stream.position();
27                if stream.expect_string(".cluster").is_ok() {
28                    return Ok(Scope::Cluster);
29                }
30                stream.set_position(saved_pos);
31            }
32            let saved_pos = stream.position();
33            // Try Cta
34            {
35                let saved_pos = stream.position();
36                if stream.expect_string(".cta").is_ok() {
37                    return Ok(Scope::Cta);
38                }
39                stream.set_position(saved_pos);
40            }
41            stream.set_position(saved_pos);
42            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
43            let expected = &[".cluster", ".cta"];
44            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
45            Err(crate::parser::unexpected_value(span, expected, found))
46        }
47    }
48
49    impl PtxParser for Sem {
50        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
51            // Try Relaxed
52            {
53                let saved_pos = stream.position();
54                if stream.expect_string(".relaxed").is_ok() {
55                    return Ok(Sem::Relaxed);
56                }
57                stream.set_position(saved_pos);
58            }
59            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
60            let expected = &[".relaxed"];
61            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
62            Err(crate::parser::unexpected_value(span, expected, found))
63        }
64    }
65
66    impl PtxParser for Space {
67        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
68            // Try SharedCluster
69            {
70                let saved_pos = stream.position();
71                if stream.expect_string(".shared::cluster").is_ok() {
72                    return Ok(Space::SharedCluster);
73                }
74                stream.set_position(saved_pos);
75            }
76            let saved_pos = stream.position();
77            // Try SharedCta
78            {
79                let saved_pos = stream.position();
80                if stream.expect_string(".shared::cta").is_ok() {
81                    return Ok(Space::SharedCta);
82                }
83                stream.set_position(saved_pos);
84            }
85            stream.set_position(saved_pos);
86            let saved_pos = stream.position();
87            // Try Shared
88            {
89                let saved_pos = stream.position();
90                if stream.expect_string(".shared").is_ok() {
91                    return Ok(Space::Shared);
92                }
93                stream.set_position(saved_pos);
94            }
95            stream.set_position(saved_pos);
96            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
97            let expected = &[".shared::cluster", ".shared::cta", ".shared"];
98            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
99            Err(crate::parser::unexpected_value(span, expected, found))
100        }
101    }
102
103    impl PtxParser for MbarrierExpectTxSemScopeSpaceB64 {
104        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
105            stream.expect_string("mbarrier")?;
106            stream.expect_string(".expect_tx")?;
107            let expect_tx = ();
108            stream.expect_complete()?;
109            let saved_pos = stream.position();
110            let sem = match Sem::parse(stream) {
111                Ok(val) => Some(val),
112                Err(_) => {
113                    stream.set_position(saved_pos);
114                    None
115                }
116            };
117            stream.expect_complete()?;
118            let saved_pos = stream.position();
119            let scope = match Scope::parse(stream) {
120                Ok(val) => Some(val),
121                Err(_) => {
122                    stream.set_position(saved_pos);
123                    None
124                }
125            };
126            stream.expect_complete()?;
127            let saved_pos = stream.position();
128            let space = match Space::parse(stream) {
129                Ok(val) => Some(val),
130                Err(_) => {
131                    stream.set_position(saved_pos);
132                    None
133                }
134            };
135            stream.expect_complete()?;
136            stream.expect_string(".b64")?;
137            let b64 = ();
138            stream.expect_complete()?;
139            let addr = AddressOperand::parse(stream)?;
140            stream.expect_complete()?;
141            stream.expect(&PtxToken::Comma)?;
142            let txcount = GeneralOperand::parse(stream)?;
143            stream.expect_complete()?;
144            stream.expect_complete()?;
145            stream.expect(&PtxToken::Semicolon)?;
146            Ok(MbarrierExpectTxSemScopeSpaceB64 {
147                expect_tx,
148                sem,
149                scope,
150                space,
151                b64,
152                addr,
153                txcount,
154            })
155        }
156    }
157
158
159}
160
ptx_parser/parser/instruction/mbarrier_expect_tx.rs

ptx_parser/parser/instruction/
mbarrier_expect_tx.rs
