ptx_parser/unparser/instruction/mbarrier_test_wait.rs
1//! Original PTX specification:
2//!
3//! mbarrier.test_wait{.sem}{.scope}{.state}.b64 waitComplete, [addr], state;
4//! mbarrier.test_wait.parity{.sem}{.scope}{.state}.b64 waitComplete, [addr], phaseParity;
5//! mbarrier.try_wait{.sem}{.scope}{.state}.b64 waitComplete, [addr], state {, suspendTimeHint};
6//! mbarrier.try_wait.parity{.sem}{.scope}{.state}.b64 waitComplete, [addr], phaseParity {, suspendTimeHint};
7//! .sem = { .acquire, .relaxed };
8//! .scope = { .cta, .cluster };
9//! .state = { .shared, .shared::cta}
10
11#![allow(unused)]
12
13use crate::lexer::PtxToken;
14use crate::unparser::{PtxUnparser, common::*};
15
16pub mod section_0 {
17 use super::*;
18 use crate::r#type::instruction::mbarrier_test_wait::section_0::*;
19
20 impl PtxUnparser for MbarrierTestWaitSemScopeStateB64 {
21 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
22 push_opcode(tokens, "mbarrier");
23 push_directive(tokens, "test_wait");
24 if let Some(sem_0) = self.sem.as_ref() {
25 match sem_0 {
26 Sem::Acquire => {
27 push_directive(tokens, "acquire");
28 }
29 Sem::Relaxed => {
30 push_directive(tokens, "relaxed");
31 }
32 }
33 }
34 if let Some(scope_1) = self.scope.as_ref() {
35 match scope_1 {
36 Scope::Cluster => {
37 push_directive(tokens, "cluster");
38 }
39 Scope::Cta => {
40 push_directive(tokens, "cta");
41 }
42 }
43 }
44 if let Some(state_2) = self.state.as_ref() {
45 match state_2 {
46 State::SharedCta => {
47 push_directive(tokens, "shared::cta");
48 }
49 State::Shared => {
50 push_directive(tokens, "shared");
51 }
52 }
53 }
54 push_directive(tokens, "b64");
55 self.waitcomplete.unparse_tokens(tokens);
56 tokens.push(PtxToken::Comma);
57 self.addr.unparse_tokens(tokens);
58 tokens.push(PtxToken::Comma);
59 self.state2.unparse_tokens(tokens);
60 tokens.push(PtxToken::Semicolon);
61 }
62 }
63
64 impl PtxUnparser for MbarrierTestWaitParitySemScopeStateB64 {
65 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
66 push_opcode(tokens, "mbarrier");
67 push_directive(tokens, "test_wait");
68 push_directive(tokens, "parity");
69 if let Some(sem_3) = self.sem.as_ref() {
70 match sem_3 {
71 Sem::Acquire => {
72 push_directive(tokens, "acquire");
73 }
74 Sem::Relaxed => {
75 push_directive(tokens, "relaxed");
76 }
77 }
78 }
79 if let Some(scope_4) = self.scope.as_ref() {
80 match scope_4 {
81 Scope::Cluster => {
82 push_directive(tokens, "cluster");
83 }
84 Scope::Cta => {
85 push_directive(tokens, "cta");
86 }
87 }
88 }
89 if let Some(state_5) = self.state.as_ref() {
90 match state_5 {
91 State::SharedCta => {
92 push_directive(tokens, "shared::cta");
93 }
94 State::Shared => {
95 push_directive(tokens, "shared");
96 }
97 }
98 }
99 push_directive(tokens, "b64");
100 self.waitcomplete.unparse_tokens(tokens);
101 tokens.push(PtxToken::Comma);
102 self.addr.unparse_tokens(tokens);
103 tokens.push(PtxToken::Comma);
104 self.phaseparity.unparse_tokens(tokens);
105 tokens.push(PtxToken::Semicolon);
106 }
107 }
108
109 impl PtxUnparser for MbarrierTryWaitSemScopeStateB64 {
110 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
111 push_opcode(tokens, "mbarrier");
112 push_directive(tokens, "try_wait");
113 if let Some(sem_6) = self.sem.as_ref() {
114 match sem_6 {
115 Sem::Acquire => {
116 push_directive(tokens, "acquire");
117 }
118 Sem::Relaxed => {
119 push_directive(tokens, "relaxed");
120 }
121 }
122 }
123 if let Some(scope_7) = self.scope.as_ref() {
124 match scope_7 {
125 Scope::Cluster => {
126 push_directive(tokens, "cluster");
127 }
128 Scope::Cta => {
129 push_directive(tokens, "cta");
130 }
131 }
132 }
133 if let Some(state_8) = self.state.as_ref() {
134 match state_8 {
135 State::SharedCta => {
136 push_directive(tokens, "shared::cta");
137 }
138 State::Shared => {
139 push_directive(tokens, "shared");
140 }
141 }
142 }
143 push_directive(tokens, "b64");
144 self.waitcomplete.unparse_tokens(tokens);
145 tokens.push(PtxToken::Comma);
146 self.addr.unparse_tokens(tokens);
147 tokens.push(PtxToken::Comma);
148 self.state2.unparse_tokens(tokens);
149 if self.suspendtimehint.is_some() { tokens.push(PtxToken::Comma); }
150 if let Some(opt_9) = self.suspendtimehint.as_ref() {
151 opt_9.unparse_tokens(tokens);
152 }
153 tokens.push(PtxToken::Semicolon);
154 }
155 }
156
157 impl PtxUnparser for MbarrierTryWaitParitySemScopeStateB64 {
158 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
159 push_opcode(tokens, "mbarrier");
160 push_directive(tokens, "try_wait");
161 push_directive(tokens, "parity");
162 if let Some(sem_10) = self.sem.as_ref() {
163 match sem_10 {
164 Sem::Acquire => {
165 push_directive(tokens, "acquire");
166 }
167 Sem::Relaxed => {
168 push_directive(tokens, "relaxed");
169 }
170 }
171 }
172 if let Some(scope_11) = self.scope.as_ref() {
173 match scope_11 {
174 Scope::Cluster => {
175 push_directive(tokens, "cluster");
176 }
177 Scope::Cta => {
178 push_directive(tokens, "cta");
179 }
180 }
181 }
182 if let Some(state_12) = self.state.as_ref() {
183 match state_12 {
184 State::SharedCta => {
185 push_directive(tokens, "shared::cta");
186 }
187 State::Shared => {
188 push_directive(tokens, "shared");
189 }
190 }
191 }
192 push_directive(tokens, "b64");
193 self.waitcomplete.unparse_tokens(tokens);
194 tokens.push(PtxToken::Comma);
195 self.addr.unparse_tokens(tokens);
196 tokens.push(PtxToken::Comma);
197 self.phaseparity.unparse_tokens(tokens);
198 if self.suspendtimehint.is_some() { tokens.push(PtxToken::Comma); }
199 if let Some(opt_13) = self.suspendtimehint.as_ref() {
200 opt_13.unparse_tokens(tokens);
201 }
202 tokens.push(PtxToken::Semicolon);
203 }
204 }
205
206}
207