ptx_parser/unparser/instruction/mbarrier_expect_tx.rs
1//! Original PTX specification:
2//!
3//! mbarrier.expect_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_expect_tx::section_0::*;
16
17 impl PtxUnparser for MbarrierExpectTxSemScopeSpaceB64 {
18 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
19 push_opcode(tokens, "mbarrier");
20 push_directive(tokens, "expect_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