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            push_opcode(tokens, "shfl");
18                    push_directive(tokens, "sync");
19                    match &self.mode {
20                            Mode::Down => {
21                                    push_directive(tokens, "down");
22                            }
23                            Mode::Bfly => {
24                                    push_directive(tokens, "bfly");
25                            }
26                            Mode::Idx => {
27                                    push_directive(tokens, "idx");
28                            }
29                            Mode::Up => {
30                                    push_directive(tokens, "up");
31                            }
32                    }
33                    push_directive(tokens, "b32");
34                    self.d.unparse_tokens(tokens);
35                    if let Some(p_0) = self.p.as_ref() {
36                        tokens.push(PtxToken::Pipe);
37                        p_0.unparse_tokens(tokens);
38                    }
39            tokens.push(PtxToken::Comma);
40                    self.a.unparse_tokens(tokens);
41            tokens.push(PtxToken::Comma);
42                    self.b.unparse_tokens(tokens);
43            tokens.push(PtxToken::Comma);
44                    self.c.unparse_tokens(tokens);
45            tokens.push(PtxToken::Comma);
46                    self.membermask.unparse_tokens(tokens);
47            tokens.push(PtxToken::Semicolon);
48        }
49    }
50
51}
52