ptx_parser/parser/instruction/
shfl_sync.rs

1//! Original PTX specification:
2//!
3//! shfl.sync.mode.b32  d{|p}, a, b, c, membermask;
4//! .mode = { .up, .down, .bfly, .idx };
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::shfl_sync::section_0::*;
15
16    // ============================================================================
17    // Generated enum parsers
18    // ============================================================================
19
20    impl PtxParser for Mode {
21        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
22            // Try Down
23            {
24                let saved_pos = stream.position();
25                if stream.expect_string(".down").is_ok() {
26                    return Ok(Mode::Down);
27                }
28                stream.set_position(saved_pos);
29            }
30            let saved_pos = stream.position();
31            // Try Bfly
32            {
33                let saved_pos = stream.position();
34                if stream.expect_string(".bfly").is_ok() {
35                    return Ok(Mode::Bfly);
36                }
37                stream.set_position(saved_pos);
38            }
39            stream.set_position(saved_pos);
40            let saved_pos = stream.position();
41            // Try Idx
42            {
43                let saved_pos = stream.position();
44                if stream.expect_string(".idx").is_ok() {
45                    return Ok(Mode::Idx);
46                }
47                stream.set_position(saved_pos);
48            }
49            stream.set_position(saved_pos);
50            let saved_pos = stream.position();
51            // Try Up
52            {
53                let saved_pos = stream.position();
54                if stream.expect_string(".up").is_ok() {
55                    return Ok(Mode::Up);
56                }
57                stream.set_position(saved_pos);
58            }
59            stream.set_position(saved_pos);
60            let span = stream
61                .peek()
62                .map(|(_, s)| s.clone())
63                .unwrap_or(Span { start: 0, end: 0 });
64            let expected = &[".down", ".bfly", ".idx", ".up"];
65            let found = stream
66                .peek()
67                .map(|(t, _)| format!("{:?}", t))
68                .unwrap_or_else(|_| "<end of input>".to_string());
69            Err(crate::parser::unexpected_value(span, expected, found))
70        }
71    }
72
73    impl PtxParser for ShflSyncModeB32 {
74        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
75            stream.expect_string("shfl")?;
76            stream.expect_string(".sync")?;
77            let sync = ();
78            stream.expect_complete()?;
79            let mode = Mode::parse(stream)?;
80            stream.expect_complete()?;
81            stream.expect_string(".b32")?;
82            let b32 = ();
83            stream.expect_complete()?;
84            let d = GeneralOperand::parse(stream)?;
85            let saved_pos = stream.position();
86            let p = if stream.consume_if(|t| matches!(t, PtxToken::Pipe)).is_some() {
87                Some(GeneralOperand::parse(stream)?)
88            } else {
89                stream.set_position(saved_pos);
90                None
91            };
92            stream.expect_complete()?;
93            stream.expect(&PtxToken::Comma)?;
94            let a = GeneralOperand::parse(stream)?;
95            stream.expect_complete()?;
96            stream.expect(&PtxToken::Comma)?;
97            let b = GeneralOperand::parse(stream)?;
98            stream.expect_complete()?;
99            stream.expect(&PtxToken::Comma)?;
100            let c = GeneralOperand::parse(stream)?;
101            stream.expect_complete()?;
102            stream.expect(&PtxToken::Comma)?;
103            let membermask = GeneralOperand::parse(stream)?;
104            stream.expect_complete()?;
105            stream.expect_complete()?;
106            stream.expect(&PtxToken::Semicolon)?;
107            Ok(ShflSyncModeB32 {
108                sync,
109                mode,
110                b32,
111                d,
112                p,
113                a,
114                b,
115                c,
116                membermask,
117            })
118        }
119    }
120}
ptx_parser/parser/instruction/shfl_sync.rs

ptx_parser/parser/instruction/
shfl_sync.rs
