ptx_parser/parser/instruction/
mbarrier_complete_tx.rs

1//! Original PTX specification:
2//!
3//! mbarrier.complete_tx{.sem}{.scope}{.space}.b64 [addr], txCount;
4//! .sem   = { .relaxed };
5//! .scope = { .cta, .cluster };
6//! .space = { .shared, .shared::cta, .shared::cluster };
7
8#![allow(unused)]
9
10use crate::parser::{
11    PtxParseError, PtxParser, PtxTokenStream, Span,
12    util::{
13        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
14        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
15    },
16};
17use crate::r#type::common::*;
18use crate::{alt, ok, seq_n};
19
20pub mod section_0 {
21    use super::*;
22    use crate::r#type::instruction::mbarrier_complete_tx::section_0::*;
23
24    // ============================================================================
25    // Generated enum parsers
26    // ============================================================================
27
28    impl PtxParser for Scope {
29        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
30            alt!(
31                map(string_p(".cluster"), |_, _span| Scope::Cluster),
32                map(string_p(".cta"), |_, _span| Scope::Cta)
33            )
34        }
35    }
36
37    impl PtxParser for Sem {
38        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
39            alt!(map(string_p(".relaxed"), |_, _span| Sem::Relaxed))
40        }
41    }
42
43    impl PtxParser for Space {
44        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
45            alt!(
46                map(string_p(".shared::cluster"), |_, _span| {
47                    Space::SharedCluster
48                }),
49                map(string_p(".shared::cta"), |_, _span| Space::SharedCta),
50                map(string_p(".shared"), |_, _span| Space::Shared)
51            )
52        }
53    }
54
55    impl PtxParser for MbarrierCompleteTxSemScopeSpaceB64 {
56        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
57            try_map(
58                seq_n!(
59                    string_p("mbarrier"),
60                    string_p(".complete_tx"),
61                    optional(Sem::parse()),
62                    optional(Scope::parse()),
63                    optional(Space::parse()),
64                    string_p(".b64"),
65                    AddressOperand::parse(),
66                    comma_p(),
67                    GeneralOperand::parse(),
68                    semicolon_p()
69                ),
70                |(_, complete_tx, sem, scope, space, b64, addr, _, txcount, _), span| {
71                    ok!(MbarrierCompleteTxSemScopeSpaceB64 {
72                        complete_tx = complete_tx,
73                        sem = sem,
74                        scope = scope,
75                        space = space,
76                        b64 = b64,
77                        addr = addr,
78                        txcount = txcount,
79
80                    })
81                },
82            )
83        }
84    }
85}
ptx_parser/parser/instruction/mbarrier_complete_tx.rs

ptx_parser/parser/instruction/
mbarrier_complete_tx.rs
