ptx_parser/unparser/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::lexer::PtxToken;
9use crate::unparser::{PtxUnparser, common::*};
10
11pub mod section_0 {
12    use super::*;
13    use crate::r#type::instruction::shfl_sync::section_0::*;
14
15    impl PtxUnparser for ShflSyncModeB32 {
16        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
17            self.unparse_tokens_mode(tokens, false);
18        }
19        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
20            push_opcode(tokens, "shfl");
21            push_directive(tokens, "sync");
22            match &self.mode {
23                Mode::Down => {
24                    push_directive(tokens, "down");
25                }
26                Mode::Bfly => {
27                    push_directive(tokens, "bfly");
28                }
29                Mode::Idx => {
30                    push_directive(tokens, "idx");
31                }
32                Mode::Up => {
33                    push_directive(tokens, "up");
34                }
35            }
36            push_directive(tokens, "b32");
37            if spaced {
38                tokens.push(PtxToken::Space);
39            }
40            self.d.unparse_tokens_mode(tokens, spaced);
41            if let Some(p_0) = self.p.as_ref() {
42                tokens.push(PtxToken::Pipe);
43                p_0.unparse_tokens_mode(tokens, spaced);
44            }
45            tokens.push(PtxToken::Comma);
46            if spaced {
47                tokens.push(PtxToken::Space);
48            }
49            self.a.unparse_tokens_mode(tokens, spaced);
50            tokens.push(PtxToken::Comma);
51            if spaced {
52                tokens.push(PtxToken::Space);
53            }
54            self.b.unparse_tokens_mode(tokens, spaced);
55            tokens.push(PtxToken::Comma);
56            if spaced {
57                tokens.push(PtxToken::Space);
58            }
59            self.c.unparse_tokens_mode(tokens, spaced);
60            tokens.push(PtxToken::Comma);
61            if spaced {
62                tokens.push(PtxToken::Space);
63            }
64            self.membermask.unparse_tokens_mode(tokens, spaced);
65            tokens.push(PtxToken::Semicolon);
66            if spaced {
67                tokens.push(PtxToken::Newline);
68            }
69        }
70    }
71}
ptx_parser/unparser/instruction/shfl_sync.rs

ptx_parser/unparser/instruction/
shfl_sync.rs
