ptx_parser/unparser/instruction/
mbarrier_arrive_drop.rs1#![allow(unused)]
13
14use crate::lexer::PtxToken;
15use crate::unparser::{PtxUnparser, common::*};
16
17pub mod section_0 {
18 use super::*;
19 use crate::r#type::instruction::mbarrier_arrive_drop::section_0::*;
20
21 impl PtxUnparser for MbarrierArriveDropSemScopeStateB64 {
22 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
23 push_opcode(tokens, "mbarrier");
24 push_directive(tokens, "arrive_drop");
25 if let Some(sem_0) = self.sem.as_ref() {
26 match sem_0 {
27 Sem::Release => {
28 push_directive(tokens, "release");
29 }
30 Sem::Relaxed => {
31 push_directive(tokens, "relaxed");
32 }
33 }
34 }
35 if let Some(scope_1) = self.scope.as_ref() {
36 match scope_1 {
37 Scope::Cluster => {
38 push_directive(tokens, "cluster");
39 }
40 Scope::Cta => {
41 push_directive(tokens, "cta");
42 }
43 }
44 }
45 if let Some(state_2) = self.state.as_ref() {
46 match state_2 {
47 State::SharedCta => {
48 push_directive(tokens, "shared::cta");
49 }
50 State::Shared => {
51 push_directive(tokens, "shared");
52 }
53 }
54 }
55 push_directive(tokens, "b64");
56 self.state2.unparse_tokens(tokens);
57 tokens.push(PtxToken::Comma);
58 self.addr.unparse_tokens(tokens);
59 if self.count.is_some() { tokens.push(PtxToken::Comma); }
60 if let Some(opt_3) = self.count.as_ref() {
61 opt_3.unparse_tokens(tokens);
62 }
63 tokens.push(PtxToken::Semicolon);
64 }
65 }
66
67 impl PtxUnparser for MbarrierArriveDropSemScopeSharedClusterB64 {
68 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
69 push_opcode(tokens, "mbarrier");
70 push_directive(tokens, "arrive_drop");
71 if let Some(sem_4) = self.sem.as_ref() {
72 match sem_4 {
73 Sem::Release => {
74 push_directive(tokens, "release");
75 }
76 Sem::Relaxed => {
77 push_directive(tokens, "relaxed");
78 }
79 }
80 }
81 if let Some(scope_5) = self.scope.as_ref() {
82 match scope_5 {
83 Scope::Cluster => {
84 push_directive(tokens, "cluster");
85 }
86 Scope::Cta => {
87 push_directive(tokens, "cta");
88 }
89 }
90 }
91 if self.shared_cluster {
92 push_directive(tokens, "shared::cluster");
93 }
94 push_directive(tokens, "b64");
95 self.operand.unparse_tokens(tokens);
96 tokens.push(PtxToken::Comma);
97 self.addr.unparse_tokens(tokens);
98 if self.count.is_some() { tokens.push(PtxToken::Comma); }
99 if let Some(opt_6) = self.count.as_ref() {
100 opt_6.unparse_tokens(tokens);
101 }
102 tokens.push(PtxToken::Semicolon);
103 }
104 }
105
106 impl PtxUnparser for MbarrierArriveDropExpectTxStateSemScopeB64 {
107 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
108 push_opcode(tokens, "mbarrier");
109 push_directive(tokens, "arrive_drop");
110 push_directive(tokens, "expect_tx");
111 if let Some(state_7) = self.state.as_ref() {
112 match state_7 {
113 State::SharedCta => {
114 push_directive(tokens, "shared::cta");
115 }
116 State::Shared => {
117 push_directive(tokens, "shared");
118 }
119 }
120 }
121 if let Some(sem_8) = self.sem.as_ref() {
122 match sem_8 {
123 Sem::Release => {
124 push_directive(tokens, "release");
125 }
126 Sem::Relaxed => {
127 push_directive(tokens, "relaxed");
128 }
129 }
130 }
131 if let Some(scope_9) = self.scope.as_ref() {
132 match scope_9 {
133 Scope::Cluster => {
134 push_directive(tokens, "cluster");
135 }
136 Scope::Cta => {
137 push_directive(tokens, "cta");
138 }
139 }
140 }
141 push_directive(tokens, "b64");
142 self.state2.unparse_tokens(tokens);
143 tokens.push(PtxToken::Comma);
144 self.addr.unparse_tokens(tokens);
145 tokens.push(PtxToken::Comma);
146 self.tx_count.unparse_tokens(tokens);
147 tokens.push(PtxToken::Semicolon);
148 }
149 }
150
151 impl PtxUnparser for MbarrierArriveDropExpectTxSharedClusterSemScopeB64 {
152 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
153 push_opcode(tokens, "mbarrier");
154 push_directive(tokens, "arrive_drop");
155 push_directive(tokens, "expect_tx");
156 if self.shared_cluster {
157 push_directive(tokens, "shared::cluster");
158 }
159 if let Some(sem_10) = self.sem.as_ref() {
160 match sem_10 {
161 Sem::Release => {
162 push_directive(tokens, "release");
163 }
164 Sem::Relaxed => {
165 push_directive(tokens, "relaxed");
166 }
167 }
168 }
169 if let Some(scope_11) = self.scope.as_ref() {
170 match scope_11 {
171 Scope::Cluster => {
172 push_directive(tokens, "cluster");
173 }
174 Scope::Cta => {
175 push_directive(tokens, "cta");
176 }
177 }
178 }
179 push_directive(tokens, "b64");
180 self.operand.unparse_tokens(tokens);
181 tokens.push(PtxToken::Comma);
182 self.addr.unparse_tokens(tokens);
183 tokens.push(PtxToken::Comma);
184 self.tx_count.unparse_tokens(tokens);
185 tokens.push(PtxToken::Semicolon);
186 }
187 }
188
189 impl PtxUnparser for MbarrierArriveDropNocompleteReleaseCtaStateB64 {
190 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
191 push_opcode(tokens, "mbarrier");
192 push_directive(tokens, "arrive_drop");
193 push_directive(tokens, "noComplete");
194 if self.release {
195 push_directive(tokens, "release");
196 }
197 if self.cta {
198 push_directive(tokens, "cta");
199 }
200 if let Some(state_12) = self.state.as_ref() {
201 match state_12 {
202 State::SharedCta => {
203 push_directive(tokens, "shared::cta");
204 }
205 State::Shared => {
206 push_directive(tokens, "shared");
207 }
208 }
209 }
210 push_directive(tokens, "b64");
211 self.state2.unparse_tokens(tokens);
212 tokens.push(PtxToken::Comma);
213 self.addr.unparse_tokens(tokens);
214 tokens.push(PtxToken::Comma);
215 self.count.unparse_tokens(tokens);
216 tokens.push(PtxToken::Semicolon);
217 }
218 }
219
220}
221