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            self.unparse_tokens_mode(tokens, false);
20        }
21        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
22            push_opcode(tokens, "mbarrier");
23            push_directive(tokens, "complete_tx");
24            if let Some(sem_0) = self.sem.as_ref() {
25                match sem_0 {
26                    Sem::Relaxed => {
27                        push_directive(tokens, "relaxed");
28                    }
29                }
30            }
31            if let Some(scope_1) = self.scope.as_ref() {
32                match scope_1 {
33                    Scope::Cluster => {
34                        push_directive(tokens, "cluster");
35                    }
36                    Scope::Cta => {
37                        push_directive(tokens, "cta");
38                    }
39                }
40            }
41            if let Some(space_2) = self.space.as_ref() {
42                match space_2 {
43                    Space::SharedCluster => {
44                        push_directive(tokens, "shared::cluster");
45                    }
46                    Space::SharedCta => {
47                        push_directive(tokens, "shared::cta");
48                    }
49                    Space::Shared => {
50                        push_directive(tokens, "shared");
51                    }
52                }
53            }
54            push_directive(tokens, "b64");
55            if spaced {
56                tokens.push(PtxToken::Space);
57            }
58            self.addr.unparse_tokens_mode(tokens, spaced);
59            tokens.push(PtxToken::Comma);
60            if spaced {
61                tokens.push(PtxToken::Space);
62            }
63            self.txcount.unparse_tokens_mode(tokens, spaced);
64            tokens.push(PtxToken::Semicolon);
65            if spaced {
66                tokens.push(PtxToken::Newline);
67            }
68        }
69    }
70}
ptx_parser/unparser/instruction/mbarrier_complete_tx.rs

ptx_parser/unparser/instruction/
mbarrier_complete_tx.rs
