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}