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
27                    .expect_string(".mbarrier::complete_tx::bytes")
28                    .is_ok()
29                {
30                    return Ok(CompletionMechanism::MbarrierCompleteTxBytes);
31                }
32                stream.set_position(saved_pos);
33            }
34            let span = stream
35                .peek()
36                .map(|(_, s)| s.clone())
37                .unwrap_or(Span { start: 0, end: 0 });
38            let expected = &[".mbarrier::complete_tx::bytes"];
39            let found = stream
40                .peek()
41                .map(|(t, _)| format!("{:?}", t))
42                .unwrap_or_else(|_| "<end of input>".to_string());
43            Err(crate::parser::unexpected_value(span, expected, found))
44        }
45    }
46
47    impl PtxParser for Space {
48        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
49            // Try SharedCta
50            {
51                let saved_pos = stream.position();
52                if stream.expect_string(".shared::cta").is_ok() {
53                    return Ok(Space::SharedCta);
54                }
55                stream.set_position(saved_pos);
56            }
57            let span = stream
58                .peek()
59                .map(|(_, s)| s.clone())
60                .unwrap_or(Span { start: 0, end: 0 });
61            let expected = &[".shared::cta"];
62            let found = stream
63                .peek()
64                .map(|(t, _)| format!("{:?}", t))
65                .unwrap_or_else(|_| "<end of input>".to_string());
66            Err(crate::parser::unexpected_value(span, expected, found))
67        }
68    }
69
70    impl PtxParser
71        for ClusterlaunchcontrolTryCancelAsyncSpaceCompletionMechanismMulticastClusterAllB128
72    {
73        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
74            stream.expect_string("clusterlaunchcontrol")?;
75            stream.expect_string(".try_cancel")?;
76            let try_cancel = ();
77            stream.expect_complete()?;
78            stream.expect_string(".async")?;
79            let async_ = ();
80            stream.expect_complete()?;
81            let saved_pos = stream.position();
82            let space = match Space::parse(stream) {
83                Ok(val) => Some(val),
84                Err(_) => {
85                    stream.set_position(saved_pos);
86                    None
87                }
88            };
89            stream.expect_complete()?;
90            let completion_mechanism = CompletionMechanism::parse(stream)?;
91            stream.expect_complete()?;
92            let saved_pos = stream.position();
93            let multicast_cluster_all = stream.expect_string(".multicast::cluster::all").is_ok();
94            if !multicast_cluster_all {
95                stream.set_position(saved_pos);
96            }
97            stream.expect_complete()?;
98            stream.expect_string(".b128")?;
99            let b128 = ();
100            stream.expect_complete()?;
101            let addr = AddressOperand::parse(stream)?;
102            stream.expect_complete()?;
103            stream.expect(&PtxToken::Comma)?;
104            let mbar = AddressOperand::parse(stream)?;
105            stream.expect_complete()?;
106            stream.expect_complete()?;
107            stream.expect(&PtxToken::Semicolon)?;
108            Ok(
109                ClusterlaunchcontrolTryCancelAsyncSpaceCompletionMechanismMulticastClusterAllB128 {
110                    try_cancel,
111                    async_,
112                    space,
113                    completion_mechanism,
114                    multicast_cluster_all,
115                    b128,
116                    addr,
117                    mbar,
118                },
119            )
120        }
121    }
122}
ptx_parser/parser/instruction/clusterlaunchcontrol_try_cancel.rs

ptx_parser/parser/instruction/
clusterlaunchcontrol_try_cancel.rs
