ptx_parser/parser/instruction/
cp_async_mbarrier_arrive.rs

1//! Original PTX specification:
2//!
3//! cp.async.mbarrier.arrive{.noinc}{.state}.b64 [addr];
4//! .state = { .shared, .shared::cta}
5
6#![allow(unused)]
7
8use crate::parser::{
9    PtxParseError, PtxParser, PtxTokenStream, Span,
10    util::{
11        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
12        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
13    },
14};
15use crate::r#type::common::*;
16use crate::{alt, ok, seq_n};
17
18pub mod section_0 {
19    use super::*;
20    use crate::r#type::instruction::cp_async_mbarrier_arrive::section_0::*;
21
22    // ============================================================================
23    // Generated enum parsers
24    // ============================================================================
25
26    impl PtxParser for State {
27        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
28            alt!(
29                map(string_p(".shared::cta"), |_, _span| State::SharedCta),
30                map(string_p(".shared"), |_, _span| State::Shared)
31            )
32        }
33    }
34
35    impl PtxParser for CpAsyncMbarrierArriveNoincStateB64 {
36        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
37            try_map(
38                seq_n!(
39                    string_p("cp"),
40                    string_p(".async"),
41                    string_p(".mbarrier"),
42                    string_p(".arrive"),
43                    map(optional(string_p(".noinc")), |value, _| value.is_some()),
44                    optional(State::parse()),
45                    string_p(".b64"),
46                    AddressOperand::parse(),
47                    semicolon_p()
48                ),
49                |(_, async_, mbarrier, arrive, noinc, state, b64, addr, _), span| {
50                    ok!(CpAsyncMbarrierArriveNoincStateB64 {
51                        async_ = async_,
52                        mbarrier = mbarrier,
53                        arrive = arrive,
54                        noinc = noinc,
55                        state = state,
56                        b64 = b64,
57                        addr = addr,
58
59                    })
60                },
61            )
62        }
63    }
64}
ptx_parser/parser/instruction/cp_async_mbarrier_arrive.rs

ptx_parser/parser/instruction/
cp_async_mbarrier_arrive.rs
