ptx_parser/parser/instruction/
tcgen05_commit.rs

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

ptx_parser/parser/instruction/
tcgen05_commit.rs
