ptx_parser/parser/instruction/
mbarrier_inval.rs1#![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 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}