ptx_parser/parser/instruction/
clusterlaunchcontrol_try_cancel.rs

1//! Original PTX specification:
2//!
3//! clusterlaunchcontrol.try_cancel.async{.space}.completion_mechanism{.multicast::cluster::all}.b128 [addr], [mbar];
4//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
5//! .space = { .shared::cta };
6
7#![allow(unused)]
8
9use crate::lexer::PtxToken;
10use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
11use crate::r#type::common::*;
12
13pub mod section_0 {
14    use super::*;
15    use crate::r#type::instruction::clusterlaunchcontrol_try_cancel::section_0::*;
16
17    // ============================================================================
18    // Generated enum parsers
19    // ============================================================================
20
21    impl PtxParser for CompletionMechanism {
22        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
23            // Try MbarrierCompleteTxBytes
24            {
25                let saved_pos = stream.position();
26                if stream.expect_string(".mbarrier::complete_tx::bytes").is_ok() {
27                    return Ok(CompletionMechanism::MbarrierCompleteTxBytes);
28                }
29                stream.set_position(saved_pos);
30            }
31            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
32            let expected = &[".mbarrier::complete_tx::bytes"];
33            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
34            Err(crate::parser::unexpected_value(span, expected, found))
35        }
36    }
37
38    impl PtxParser for Space {
39        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
40            // Try SharedCta
41            {
42                let saved_pos = stream.position();
43                if stream.expect_string(".shared::cta").is_ok() {
44                    return Ok(Space::SharedCta);
45                }
46                stream.set_position(saved_pos);
47            }
48            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
49            let expected = &[".shared::cta"];
50            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
51            Err(crate::parser::unexpected_value(span, expected, found))
52        }
53    }
54
55    impl PtxParser for ClusterlaunchcontrolTryCancelAsyncSpaceCompletionMechanismMulticastClusterAllB128 {
56        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
57            stream.expect_string("clusterlaunchcontrol")?;
58            stream.expect_string(".try_cancel")?;
59            let try_cancel = ();
60            stream.expect_complete()?;
61            stream.expect_string(".async")?;
62            let async_ = ();
63            stream.expect_complete()?;
64            let saved_pos = stream.position();
65            let space = match Space::parse(stream) {
66                Ok(val) => Some(val),
67                Err(_) => {
68                    stream.set_position(saved_pos);
69                    None
70                }
71            };
72            stream.expect_complete()?;
73            let completion_mechanism = CompletionMechanism::parse(stream)?;
74            stream.expect_complete()?;
75            let saved_pos = stream.position();
76            let multicast_cluster_all = stream.expect_string(".multicast::cluster::all").is_ok();
77            if !multicast_cluster_all {
78                stream.set_position(saved_pos);
79            }
80            stream.expect_complete()?;
81            stream.expect_string(".b128")?;
82            let b128 = ();
83            stream.expect_complete()?;
84            let addr = AddressOperand::parse(stream)?;
85            stream.expect_complete()?;
86            stream.expect(&PtxToken::Comma)?;
87            let mbar = AddressOperand::parse(stream)?;
88            stream.expect_complete()?;
89            stream.expect_complete()?;
90            stream.expect(&PtxToken::Semicolon)?;
91            Ok(ClusterlaunchcontrolTryCancelAsyncSpaceCompletionMechanismMulticastClusterAllB128 {
92                try_cancel,
93                async_,
94                space,
95                completion_mechanism,
96                multicast_cluster_all,
97                b128,
98                addr,
99                mbar,
100            })
101        }
102    }
103
104
105}
106
ptx_parser/parser/instruction/clusterlaunchcontrol_try_cancel.rs

ptx_parser/parser/instruction/
clusterlaunchcontrol_try_cancel.rs
