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::parser::{
12    PtxParseError, PtxParser, PtxTokenStream, Span,
13    util::{
14        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
15        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
16    },
17};
18use crate::r#type::common::*;
19use crate::{alt, ok, seq_n};
20
21pub mod section_0 {
22    use super::*;
23    use crate::r#type::instruction::tcgen05_commit::section_0::*;
24
25    // ============================================================================
26    // Generated enum parsers
27    // ============================================================================
28
29    impl PtxParser for CompletionMechanism {
30        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
31            alt!(map(string_p(".mbarrier::arrive::one"), |_, _span| {
32                CompletionMechanism::MbarrierArriveOne
33            }))
34        }
35    }
36
37    impl PtxParser for CtaGroup {
38        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
39            alt!(
40                map(string_p(".cta_group::1"), |_, _span| CtaGroup::CtaGroup1),
41                map(string_p(".cta_group::2"), |_, _span| CtaGroup::CtaGroup2)
42            )
43        }
44    }
45
46    impl PtxParser for Multicast {
47        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
48            alt!(map(string_p(".multicast::cluster"), |_, _span| {
49                Multicast::MulticastCluster
50            }))
51        }
52    }
53
54    impl PtxParser for Tcgen05CommitCtaGroupCompletionMechanismSharedClusterMulticastB64 {
55        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
56            try_map(
57                seq_n!(
58                    string_p("tcgen05"),
59                    string_p(".commit"),
60                    CtaGroup::parse(),
61                    CompletionMechanism::parse(),
62                    map(optional(string_p(".shared::cluster")), |value, _| value
63                        .is_some()),
64                    optional(Multicast::parse()),
65                    string_p(".b64"),
66                    AddressOperand::parse(),
67                    map(
68                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
69                        |value, _| value.map(|(_, operand)| operand)
70                    ),
71                    semicolon_p()
72                ),
73                |(
74                    _,
75                    commit,
76                    cta_group,
77                    completion_mechanism,
78                    shared_cluster,
79                    multicast,
80                    b64,
81                    mbar,
82                    ctamask,
83                    _,
84                ),
85                 span| {
86                    ok!(Tcgen05CommitCtaGroupCompletionMechanismSharedClusterMulticastB64 {
87                        commit = commit,
88                        cta_group = cta_group,
89                        completion_mechanism = completion_mechanism,
90                        shared_cluster = shared_cluster,
91                        multicast = multicast,
92                        b64 = b64,
93                        mbar = mbar,
94                        ctamask = ctamask,
95
96                    })
97                },
98            )
99        }
100    }
101}
ptx_parser/parser/instruction/tcgen05_commit.rs

ptx_parser/parser/instruction/
tcgen05_commit.rs
