Skip to main content

ptx_parser/parser/instruction/
mbarrier_inval.rs

1//! Original PTX specification:
2//!
3//! mbarrier.inval{.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::mbarrier_inval::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 MbarrierInvalStateB64 {
36        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
37            try_map(
38                seq_n!(
39                    string_p("mbarrier"),
40                    string_p(".inval"),
41                    optional(State::parse()),
42                    string_p(".b64"),
43                    AddressOperand::parse(),
44                    semicolon_p()
45                ),
46                |(_, inval, state, b64, addr, _), span| {
47                    ok!(MbarrierInvalStateB64 {
48                        inval = inval,
49                        state = state,
50                        b64 = b64,
51                        addr = addr,
52
53                    })
54                },
55            )
56        }
57    }
58}