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
34                .peek()
35                .map(|(_, s)| s.clone())
36                .unwrap_or(Span { start: 0, end: 0 });
37            let expected = &[".mbarrier::arrive::one"];
38            let found = stream
39                .peek()
40                .map(|(t, _)| format!("{:?}", t))
41                .unwrap_or_else(|_| "<end of input>".to_string());
42            Err(crate::parser::unexpected_value(span, expected, found))
43        }
44    }
45
46    impl PtxParser for CtaGroup {
47        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
48            // Try CtaGroup1
49            {
50                let saved_pos = stream.position();
51                if stream.expect_string(".cta_group::1").is_ok() {
52                    return Ok(CtaGroup::CtaGroup1);
53                }
54                stream.set_position(saved_pos);
55            }
56            let saved_pos = stream.position();
57            // Try CtaGroup2
58            {
59                let saved_pos = stream.position();
60                if stream.expect_string(".cta_group::2").is_ok() {
61                    return Ok(CtaGroup::CtaGroup2);
62                }
63                stream.set_position(saved_pos);
64            }
65            stream.set_position(saved_pos);
66            let span = stream
67                .peek()
68                .map(|(_, s)| s.clone())
69                .unwrap_or(Span { start: 0, end: 0 });
70            let expected = &[".cta_group::1", ".cta_group::2"];
71            let found = stream
72                .peek()
73                .map(|(t, _)| format!("{:?}", t))
74                .unwrap_or_else(|_| "<end of input>".to_string());
75            Err(crate::parser::unexpected_value(span, expected, found))
76        }
77    }
78
79    impl PtxParser for Multicast {
80        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
81            // Try MulticastCluster
82            {
83                let saved_pos = stream.position();
84                if stream.expect_string(".multicast::cluster").is_ok() {
85                    return Ok(Multicast::MulticastCluster);
86                }
87                stream.set_position(saved_pos);
88            }
89            let span = stream
90                .peek()
91                .map(|(_, s)| s.clone())
92                .unwrap_or(Span { start: 0, end: 0 });
93            let expected = &[".multicast::cluster"];
94            let found = stream
95                .peek()
96                .map(|(t, _)| format!("{:?}", t))
97                .unwrap_or_else(|_| "<end of input>".to_string());
98            Err(crate::parser::unexpected_value(span, expected, found))
99        }
100    }
101
102    impl PtxParser for Tcgen05CommitCtaGroupCompletionMechanismSharedClusterMulticastB64 {
103        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
104            stream.expect_string("tcgen05")?;
105            stream.expect_string(".commit")?;
106            let commit = ();
107            stream.expect_complete()?;
108            let cta_group = CtaGroup::parse(stream)?;
109            stream.expect_complete()?;
110            let completion_mechanism = CompletionMechanism::parse(stream)?;
111            stream.expect_complete()?;
112            let saved_pos = stream.position();
113            let shared_cluster = stream.expect_string(".shared::cluster").is_ok();
114            if !shared_cluster {
115                stream.set_position(saved_pos);
116            }
117            stream.expect_complete()?;
118            let saved_pos = stream.position();
119            let multicast = match Multicast::parse(stream) {
120                Ok(val) => Some(val),
121                Err(_) => {
122                    stream.set_position(saved_pos);
123                    None
124                }
125            };
126            stream.expect_complete()?;
127            stream.expect_string(".b64")?;
128            let b64 = ();
129            stream.expect_complete()?;
130            let mbar = AddressOperand::parse(stream)?;
131            stream.expect_complete()?;
132            let saved_pos = stream.position();
133            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
134            if !has_comma {
135                stream.set_position(saved_pos);
136            }
137            let saved_pos = stream.position();
138            let ctamask = match GeneralOperand::parse(stream) {
139                Ok(val) => Some(val),
140                Err(_) => {
141                    stream.set_position(saved_pos);
142                    None
143                }
144            };
145            stream.expect_complete()?;
146            stream.expect_complete()?;
147            stream.expect(&PtxToken::Semicolon)?;
148            Ok(
149                Tcgen05CommitCtaGroupCompletionMechanismSharedClusterMulticastB64 {
150                    commit,
151                    cta_group,
152                    completion_mechanism,
153                    shared_cluster,
154                    multicast,
155                    b64,
156                    mbar,
157                    ctamask,
158                },
159            )
160        }
161    }
162}
ptx_parser/parser/instruction/tcgen05_commit.rs

ptx_parser/parser/instruction/
tcgen05_commit.rs
