ptx_parser/unparser/instruction/
mbarrier_init.rs

1//! Original PTX specification:
2//!
3//! mbarrier.init{.state}.b64 [addr], count;
4//! .state = { .shared, .shared::cta}
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::mbarrier_init::section_0::*;
14
15    impl PtxUnparser for MbarrierInitStateB64 {
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, "mbarrier");
21            push_directive(tokens, "init");
22            if let Some(state_0) = self.state.as_ref() {
23                match state_0 {
24                    State::SharedCta => {
25                        push_directive(tokens, "shared::cta");
26                    }
27                    State::Shared => {
28                        push_directive(tokens, "shared");
29                    }
30                }
31            }
32            push_directive(tokens, "b64");
33            if spaced {
34                tokens.push(PtxToken::Space);
35            }
36            self.addr.unparse_tokens_mode(tokens, spaced);
37            tokens.push(PtxToken::Comma);
38            if spaced {
39                tokens.push(PtxToken::Space);
40            }
41            self.count.unparse_tokens_mode(tokens, spaced);
42            tokens.push(PtxToken::Semicolon);
43            if spaced {
44                tokens.push(PtxToken::Newline);
45            }
46        }
47    }
48}
ptx_parser/unparser/instruction/mbarrier_init.rs

ptx_parser/unparser/instruction/
mbarrier_init.rs
