ptx_parser/parser/instruction/
tcgen05_shift.rs

1//! Original PTX specification:
2//!
3//! tcgen05.shift.cta_group.down  [taddr];
4//! .cta_group = { .cta_group::1, .cta_group::2 }
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::tcgen05_shift::section_0::*;
15
16    // ============================================================================
17    // Generated enum parsers
18    // ============================================================================
19
20    impl PtxParser for CtaGroup {
21        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
22            // Try CtaGroup1
23            {
24                let saved_pos = stream.position();
25                if stream.expect_string(".cta_group::1").is_ok() {
26                    return Ok(CtaGroup::CtaGroup1);
27                }
28                stream.set_position(saved_pos);
29            }
30            let saved_pos = stream.position();
31            // Try CtaGroup2
32            {
33                let saved_pos = stream.position();
34                if stream.expect_string(".cta_group::2").is_ok() {
35                    return Ok(CtaGroup::CtaGroup2);
36                }
37                stream.set_position(saved_pos);
38            }
39            stream.set_position(saved_pos);
40            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
41            let expected = &[".cta_group::1", ".cta_group::2"];
42            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
43            Err(crate::parser::unexpected_value(span, expected, found))
44        }
45    }
46
47    impl PtxParser for Tcgen05ShiftCtaGroupDown {
48        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
49            stream.expect_string("tcgen05")?;
50            stream.expect_string(".shift")?;
51            let shift = ();
52            stream.expect_complete()?;
53            let cta_group = CtaGroup::parse(stream)?;
54            stream.expect_complete()?;
55            stream.expect_string(".down")?;
56            let down = ();
57            stream.expect_complete()?;
58            let taddr = AddressOperand::parse(stream)?;
59            stream.expect_complete()?;
60            stream.expect_complete()?;
61            stream.expect(&PtxToken::Semicolon)?;
62            Ok(Tcgen05ShiftCtaGroupDown {
63                shift,
64                cta_group,
65                down,
66                taddr,
67            })
68        }
69    }
70
71
72}
73
ptx_parser/parser/instruction/tcgen05_shift.rs

ptx_parser/parser/instruction/
tcgen05_shift.rs
