ptx_parser/parser/instruction/
shfl_sync.rs

1//! Original PTX specification:
2//!
3//! shfl.sync.mode.b32  d{|p}, a, b, c, membermask;
4//! .mode = { .up, .down, .bfly, .idx };
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::shfl_sync::section_0::*;
21
22    // ============================================================================
23    // Generated enum parsers
24    // ============================================================================
25
26    impl PtxParser for Mode {
27        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
28            alt!(
29                map(string_p(".down"), |_, _span| Mode::Down),
30                map(string_p(".bfly"), |_, _span| Mode::Bfly),
31                map(string_p(".idx"), |_, _span| Mode::Idx),
32                map(string_p(".up"), |_, _span| Mode::Up)
33            )
34        }
35    }
36
37    impl PtxParser for ShflSyncModeB32 {
38        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
39            try_map(
40                seq_n!(
41                    string_p("shfl"),
42                    string_p(".sync"),
43                    Mode::parse(),
44                    string_p(".b32"),
45                    GeneralOperand::parse(),
46                    map(
47                        optional(seq_n!(pipe_p(), GeneralOperand::parse())),
48                        |value, _| value.map(|(_, operand)| operand)
49                    ),
50                    comma_p(),
51                    GeneralOperand::parse(),
52                    comma_p(),
53                    GeneralOperand::parse(),
54                    comma_p(),
55                    GeneralOperand::parse(),
56                    comma_p(),
57                    GeneralOperand::parse(),
58                    semicolon_p()
59                ),
60                |(_, sync, mode, b32, d, p, _, a, _, b, _, c, _, membermask, _), span| {
61                    ok!(ShflSyncModeB32 {
62                        sync = sync,
63                        mode = mode,
64                        b32 = b32,
65                        d = d,
66                        p = p,
67                        a = a,
68                        b = b,
69                        c = c,
70                        membermask = membermask,
71
72                    })
73                },
74            )
75        }
76    }
77}
ptx_parser/parser/instruction/shfl_sync.rs

ptx_parser/parser/instruction/
shfl_sync.rs
