ptx_parser/parser/instruction/
mbarrier_complete_tx.rs

1//! Original PTX specification:
2//!
3//! mbarrier.complete_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_complete_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
43                .peek()
44                .map(|(_, s)| s.clone())
45                .unwrap_or(Span { start: 0, end: 0 });
46            let expected = &[".cluster", ".cta"];
47            let found = stream
48                .peek()
49                .map(|(t, _)| format!("{:?}", t))
50                .unwrap_or_else(|_| "<end of input>".to_string());
51            Err(crate::parser::unexpected_value(span, expected, found))
52        }
53    }
54
55    impl PtxParser for Sem {
56        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
57            // Try Relaxed
58            {
59                let saved_pos = stream.position();
60                if stream.expect_string(".relaxed").is_ok() {
61                    return Ok(Sem::Relaxed);
62                }
63                stream.set_position(saved_pos);
64            }
65            let span = stream
66                .peek()
67                .map(|(_, s)| s.clone())
68                .unwrap_or(Span { start: 0, end: 0 });
69            let expected = &[".relaxed"];
70            let found = stream
71                .peek()
72                .map(|(t, _)| format!("{:?}", t))
73                .unwrap_or_else(|_| "<end of input>".to_string());
74            Err(crate::parser::unexpected_value(span, expected, found))
75        }
76    }
77
78    impl PtxParser for Space {
79        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
80            // Try SharedCluster
81            {
82                let saved_pos = stream.position();
83                if stream.expect_string(".shared::cluster").is_ok() {
84                    return Ok(Space::SharedCluster);
85                }
86                stream.set_position(saved_pos);
87            }
88            let saved_pos = stream.position();
89            // Try SharedCta
90            {
91                let saved_pos = stream.position();
92                if stream.expect_string(".shared::cta").is_ok() {
93                    return Ok(Space::SharedCta);
94                }
95                stream.set_position(saved_pos);
96            }
97            stream.set_position(saved_pos);
98            let saved_pos = stream.position();
99            // Try Shared
100            {
101                let saved_pos = stream.position();
102                if stream.expect_string(".shared").is_ok() {
103                    return Ok(Space::Shared);
104                }
105                stream.set_position(saved_pos);
106            }
107            stream.set_position(saved_pos);
108            let span = stream
109                .peek()
110                .map(|(_, s)| s.clone())
111                .unwrap_or(Span { start: 0, end: 0 });
112            let expected = &[".shared::cluster", ".shared::cta", ".shared"];
113            let found = stream
114                .peek()
115                .map(|(t, _)| format!("{:?}", t))
116                .unwrap_or_else(|_| "<end of input>".to_string());
117            Err(crate::parser::unexpected_value(span, expected, found))
118        }
119    }
120
121    impl PtxParser for MbarrierCompleteTxSemScopeSpaceB64 {
122        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
123            stream.expect_string("mbarrier")?;
124            stream.expect_string(".complete_tx")?;
125            let complete_tx = ();
126            stream.expect_complete()?;
127            let saved_pos = stream.position();
128            let sem = match Sem::parse(stream) {
129                Ok(val) => Some(val),
130                Err(_) => {
131                    stream.set_position(saved_pos);
132                    None
133                }
134            };
135            stream.expect_complete()?;
136            let saved_pos = stream.position();
137            let scope = match Scope::parse(stream) {
138                Ok(val) => Some(val),
139                Err(_) => {
140                    stream.set_position(saved_pos);
141                    None
142                }
143            };
144            stream.expect_complete()?;
145            let saved_pos = stream.position();
146            let space = match Space::parse(stream) {
147                Ok(val) => Some(val),
148                Err(_) => {
149                    stream.set_position(saved_pos);
150                    None
151                }
152            };
153            stream.expect_complete()?;
154            stream.expect_string(".b64")?;
155            let b64 = ();
156            stream.expect_complete()?;
157            let addr = AddressOperand::parse(stream)?;
158            stream.expect_complete()?;
159            stream.expect(&PtxToken::Comma)?;
160            let txcount = GeneralOperand::parse(stream)?;
161            stream.expect_complete()?;
162            stream.expect_complete()?;
163            stream.expect(&PtxToken::Semicolon)?;
164            Ok(MbarrierCompleteTxSemScopeSpaceB64 {
165                complete_tx,
166                sem,
167                scope,
168                space,
169                b64,
170                addr,
171                txcount,
172            })
173        }
174    }
175}
ptx_parser/parser/instruction/mbarrier_complete_tx.rs

ptx_parser/parser/instruction/
mbarrier_complete_tx.rs
