ptx_parser/unparser/instruction/
cp_async_mbarrier_arrive.rs

1//! Original PTX specification:
2//!
3//! cp.async.mbarrier.arrive{.noinc}{.state}.b64 [addr];
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::cp_async_mbarrier_arrive::section_0::*;
14
15    impl PtxUnparser for CpAsyncMbarrierArriveNoincStateB64 {
16        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
17            push_opcode(tokens, "cp");
18                    push_directive(tokens, "async");
19                    push_directive(tokens, "mbarrier");
20                    push_directive(tokens, "arrive");
21                    if self.noinc {
22                            push_directive(tokens, "noinc");
23                    }
24                    if let Some(state_0) = self.state.as_ref() {
25                            match state_0 {
26                                    State::SharedCta => {
27                                            push_directive(tokens, "shared::cta");
28                                    }
29                                    State::Shared => {
30                                            push_directive(tokens, "shared");
31                                    }
32                            }
33                    }
34                    push_directive(tokens, "b64");
35                    self.addr.unparse_tokens(tokens);
36            tokens.push(PtxToken::Semicolon);
37        }
38    }
39
40}
41