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::parser::{
10    PtxParseError, PtxParser, PtxTokenStream, Span,
11    util::{
12        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
13        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
14    },
15};
16use crate::r#type::common::*;
17use crate::{alt, ok, seq_n};
18
19pub mod section_0 {
20    use super::*;
21    use crate::r#type::instruction::clusterlaunchcontrol_try_cancel::section_0::*;
22
23    // ============================================================================
24    // Generated enum parsers
25    // ============================================================================
26
27    impl PtxParser for CompletionMechanism {
28        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
29            alt!(map(
30                string_p(".mbarrier::complete_tx::bytes"),
31                |_, _span| CompletionMechanism::MbarrierCompleteTxBytes
32            ))
33        }
34    }
35
36    impl PtxParser for Space {
37        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
38            alt!(map(string_p(".shared::cta"), |_, _span| Space::SharedCta))
39        }
40    }
41
42    impl PtxParser
43        for ClusterlaunchcontrolTryCancelAsyncSpaceCompletionMechanismMulticastClusterAllB128
44    {
45        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
46            try_map(
47                seq_n!(
48                    string_p("clusterlaunchcontrol"),
49                    string_p(".try_cancel"),
50                    string_p(".async"),
51                    optional(Space::parse()),
52                    CompletionMechanism::parse(),
53                    map(
54                        optional(string_p(".multicast::cluster::all")),
55                        |value, _| value.is_some()
56                    ),
57                    string_p(".b128"),
58                    AddressOperand::parse(),
59                    comma_p(),
60                    AddressOperand::parse(),
61                    semicolon_p()
62                ),
63                |(
64                    _,
65                    try_cancel,
66                    async_,
67                    space,
68                    completion_mechanism,
69                    multicast_cluster_all,
70                    b128,
71                    addr,
72                    _,
73                    mbar,
74                    _,
75                ),
76                 span| {
77                    ok!(ClusterlaunchcontrolTryCancelAsyncSpaceCompletionMechanismMulticastClusterAllB128 {
78                        try_cancel = try_cancel,
79                        async_ = async_,
80                        space = space,
81                        completion_mechanism = completion_mechanism,
82                        multicast_cluster_all = multicast_cluster_all,
83                        b128 = b128,
84                        addr = addr,
85                        mbar = mbar,
86
87                    })
88                },
89            )
90        }
91    }
92}
ptx_parser/parser/instruction/clusterlaunchcontrol_try_cancel.rs

ptx_parser/parser/instruction/
clusterlaunchcontrol_try_cancel.rs
