ptx_parser/parser/instruction/
mbarrier_init.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_init::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 MbarrierInitStateB64 {
36 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
37 try_map(
38 seq_n!(
39 string_p("mbarrier"),
40 string_p(".init"),
41 optional(State::parse()),
42 string_p(".b64"),
43 AddressOperand::parse(),
44 comma_p(),
45 GeneralOperand::parse(),
46 semicolon_p()
47 ),
48 |(_, init, state, b64, addr, _, count, _), span| {
49 ok!(MbarrierInitStateB64 {
50 init = init,
51 state = state,
52 b64 = b64,
53 addr = addr,
54 count = count,
55
56 })
57 },
58 )
59 }
60 }
61}