ptx_parser/unparser/instruction/
mbarrier_complete_tx.rs

1//! Original PTX specification:
2//!
3//! mbarrier.complete_tx{.sem}{.scope}{.space}.b64 [addr], txCount;
4//! .sem   = { .relaxed };
5//! .scope = { .cta, .cluster };
6//! .space = { .shared, .shared::cta, .shared::cluster };
7
8#![allow(unused)]
9
10use crate::lexer::PtxToken;
11use crate::unparser::{PtxUnparser, common::*};
12
13pub mod section_0 {
14    use super::*;
15    use crate::r#type::instruction::mbarrier_complete_tx::section_0::*;
16
17    impl PtxUnparser for MbarrierCompleteTxSemScopeSpaceB64 {
18        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
19            push_opcode(tokens, "mbarrier");
20                    push_directive(tokens, "complete_tx");
21                    if let Some(sem_0) = self.sem.as_ref() {
22                            match sem_0 {
23                                    Sem::Relaxed => {
24                                            push_directive(tokens, "relaxed");
25                                    }
26                            }
27                    }
28                    if let Some(scope_1) = self.scope.as_ref() {
29                            match scope_1 {
30                                    Scope::Cluster => {
31                                            push_directive(tokens, "cluster");
32                                    }
33                                    Scope::Cta => {
34                                            push_directive(tokens, "cta");
35                                    }
36                            }
37                    }
38                    if let Some(space_2) = self.space.as_ref() {
39                            match space_2 {
40                                    Space::SharedCluster => {
41                                            push_directive(tokens, "shared::cluster");
42                                    }
43                                    Space::SharedCta => {
44                                            push_directive(tokens, "shared::cta");
45                                    }
46                                    Space::Shared => {
47                                            push_directive(tokens, "shared");
48                                    }
49                            }
50                    }
51                    push_directive(tokens, "b64");
52                    self.addr.unparse_tokens(tokens);
53            tokens.push(PtxToken::Comma);
54                    self.txcount.unparse_tokens(tokens);
55            tokens.push(PtxToken::Semicolon);
56        }
57    }
58
59}
60