ptx_parser/unparser/instruction/red_async.rs
1//! Original PTX specification:
2//!
3//! // Increment and Decrement reductions
4//! red.async.sem.scope{.ss}.completion_mechanism.op.type [a], b, [mbar];
5//! .sem = { .relaxed };
6//! .scope = { .cluster };
7//! .ss = { .shared::cluster };
8//! .op = { .inc, .dec };
9//! .type = { .u32 };
10//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
11//! ------------------------------------------------------------------
12//! // MIN and MAX reductions
13//! red.async.sem.scope{.ss}.completion_mechanism.op.type [a], b, [mbar];
14//! .sem = { .relaxed };
15//! .scope = { .cluster };
16//! .ss = { .shared::cluster };
17//! .op = { .min, .max };
18//! .type = { .u32, .s32 };
19//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
20//! ------------------------------------------------------------------
21//! // Bitwise AND, OR and XOR reductions
22//! red.async.sem.scope{.ss}.completion_mechanism.op.type [a], b, [mbar];
23//! .sem = { .relaxed };
24//! .scope = { .cluster };
25//! .ss = { .shared::cluster };
26//! .op = { .and, .or, .xor };
27//! .type = { .b32 };
28//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
29//! ------------------------------------------------------------------
30//! // ADD reductions
31//! red.async.sem.scope{.ss}.completion_mechanism.add.type [a], b, [mbar];
32//! .sem = { .relaxed };
33//! .scope = { .cluster };
34//! .ss = { .shared::cluster };
35//! .type = { .u32, .s32, .u64 };
36//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
37//! ----------------------------------------------------
38//! // Alternate floating point type:
39//! red.async{.mmio}.sem.scope{.ss}.add.type [a], b;
40//! .sem = { .release };
41//! .scope = { .gpu, .cluster };
42//! .ss = { .global };
43//! .type = { .u32, .s32, .u64, .s64 };
44
45#![allow(unused)]
46
47use crate::lexer::PtxToken;
48use crate::unparser::{PtxUnparser, common::*};
49
50pub mod section_0 {
51 use super::*;
52 use crate::r#type::instruction::red_async::section_0::*;
53
54 impl PtxUnparser for RedAsyncSemScopeSsCompletionMechanismOpType {
55 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
56 push_opcode(tokens, "red");
57 push_directive(tokens, "async");
58 match &self.sem {
59 Sem::Relaxed => {
60 push_directive(tokens, "relaxed");
61 }
62 }
63 match &self.scope {
64 Scope::Cluster => {
65 push_directive(tokens, "cluster");
66 }
67 }
68 if let Some(ss_0) = self.ss.as_ref() {
69 match ss_0 {
70 Ss::SharedCluster => {
71 push_directive(tokens, "shared::cluster");
72 }
73 }
74 }
75 match &self.completion_mechanism {
76 CompletionMechanism::MbarrierCompleteTxBytes => {
77 push_directive(tokens, "mbarrier::complete_tx::bytes");
78 }
79 }
80 match &self.op {
81 Op::Inc => {
82 push_directive(tokens, "inc");
83 }
84 Op::Dec => {
85 push_directive(tokens, "dec");
86 }
87 }
88 match &self.type_ {
89 Type::U32 => {
90 push_directive(tokens, "u32");
91 }
92 }
93 self.a.unparse_tokens(tokens);
94 tokens.push(PtxToken::Comma);
95 self.b.unparse_tokens(tokens);
96 tokens.push(PtxToken::Comma);
97 self.mbar.unparse_tokens(tokens);
98 tokens.push(PtxToken::Semicolon);
99 }
100 }
101
102}
103
104pub mod section_1 {
105 use super::*;
106 use crate::r#type::instruction::red_async::section_1::*;
107
108 impl PtxUnparser for RedAsyncSemScopeSsCompletionMechanismOpType1 {
109 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
110 push_opcode(tokens, "red");
111 push_directive(tokens, "async");
112 match &self.sem {
113 Sem::Relaxed => {
114 push_directive(tokens, "relaxed");
115 }
116 }
117 match &self.scope {
118 Scope::Cluster => {
119 push_directive(tokens, "cluster");
120 }
121 }
122 if let Some(ss_1) = self.ss.as_ref() {
123 match ss_1 {
124 Ss::SharedCluster => {
125 push_directive(tokens, "shared::cluster");
126 }
127 }
128 }
129 match &self.completion_mechanism {
130 CompletionMechanism::MbarrierCompleteTxBytes => {
131 push_directive(tokens, "mbarrier::complete_tx::bytes");
132 }
133 }
134 match &self.op {
135 Op::Min => {
136 push_directive(tokens, "min");
137 }
138 Op::Max => {
139 push_directive(tokens, "max");
140 }
141 }
142 match &self.type_ {
143 Type::U32 => {
144 push_directive(tokens, "u32");
145 }
146 Type::S32 => {
147 push_directive(tokens, "s32");
148 }
149 }
150 self.a.unparse_tokens(tokens);
151 tokens.push(PtxToken::Comma);
152 self.b.unparse_tokens(tokens);
153 tokens.push(PtxToken::Comma);
154 self.mbar.unparse_tokens(tokens);
155 tokens.push(PtxToken::Semicolon);
156 }
157 }
158
159}
160
161pub mod section_2 {
162 use super::*;
163 use crate::r#type::instruction::red_async::section_2::*;
164
165 impl PtxUnparser for RedAsyncSemScopeSsCompletionMechanismOpType2 {
166 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
167 push_opcode(tokens, "red");
168 push_directive(tokens, "async");
169 match &self.sem {
170 Sem::Relaxed => {
171 push_directive(tokens, "relaxed");
172 }
173 }
174 match &self.scope {
175 Scope::Cluster => {
176 push_directive(tokens, "cluster");
177 }
178 }
179 if let Some(ss_2) = self.ss.as_ref() {
180 match ss_2 {
181 Ss::SharedCluster => {
182 push_directive(tokens, "shared::cluster");
183 }
184 }
185 }
186 match &self.completion_mechanism {
187 CompletionMechanism::MbarrierCompleteTxBytes => {
188 push_directive(tokens, "mbarrier::complete_tx::bytes");
189 }
190 }
191 match &self.op {
192 Op::And => {
193 push_directive(tokens, "and");
194 }
195 Op::Xor => {
196 push_directive(tokens, "xor");
197 }
198 Op::Or => {
199 push_directive(tokens, "or");
200 }
201 }
202 match &self.type_ {
203 Type::B32 => {
204 push_directive(tokens, "b32");
205 }
206 }
207 self.a.unparse_tokens(tokens);
208 tokens.push(PtxToken::Comma);
209 self.b.unparse_tokens(tokens);
210 tokens.push(PtxToken::Comma);
211 self.mbar.unparse_tokens(tokens);
212 tokens.push(PtxToken::Semicolon);
213 }
214 }
215
216}
217
218pub mod section_3 {
219 use super::*;
220 use crate::r#type::instruction::red_async::section_3::*;
221
222 impl PtxUnparser for RedAsyncSemScopeSsCompletionMechanismAddType {
223 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
224 push_opcode(tokens, "red");
225 push_directive(tokens, "async");
226 match &self.sem {
227 Sem::Relaxed => {
228 push_directive(tokens, "relaxed");
229 }
230 }
231 match &self.scope {
232 Scope::Cluster => {
233 push_directive(tokens, "cluster");
234 }
235 }
236 if let Some(ss_3) = self.ss.as_ref() {
237 match ss_3 {
238 Ss::SharedCluster => {
239 push_directive(tokens, "shared::cluster");
240 }
241 }
242 }
243 match &self.completion_mechanism {
244 CompletionMechanism::MbarrierCompleteTxBytes => {
245 push_directive(tokens, "mbarrier::complete_tx::bytes");
246 }
247 }
248 push_directive(tokens, "add");
249 match &self.type_ {
250 Type::U32 => {
251 push_directive(tokens, "u32");
252 }
253 Type::S32 => {
254 push_directive(tokens, "s32");
255 }
256 Type::U64 => {
257 push_directive(tokens, "u64");
258 }
259 }
260 self.a.unparse_tokens(tokens);
261 tokens.push(PtxToken::Comma);
262 self.b.unparse_tokens(tokens);
263 tokens.push(PtxToken::Comma);
264 self.mbar.unparse_tokens(tokens);
265 tokens.push(PtxToken::Semicolon);
266 }
267 }
268
269}
270
271pub mod section_4 {
272 use super::*;
273 use crate::r#type::instruction::red_async::section_4::*;
274
275 impl PtxUnparser for RedAsyncMmioSemScopeSsAddType {
276 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
277 push_opcode(tokens, "red");
278 push_directive(tokens, "async");
279 if self.mmio {
280 push_directive(tokens, "mmio");
281 }
282 match &self.sem {
283 Sem::Release => {
284 push_directive(tokens, "release");
285 }
286 }
287 match &self.scope {
288 Scope::Cluster => {
289 push_directive(tokens, "cluster");
290 }
291 Scope::Gpu => {
292 push_directive(tokens, "gpu");
293 }
294 }
295 if let Some(ss_4) = self.ss.as_ref() {
296 match ss_4 {
297 Ss::Global => {
298 push_directive(tokens, "global");
299 }
300 }
301 }
302 push_directive(tokens, "add");
303 match &self.type_ {
304 Type::U32 => {
305 push_directive(tokens, "u32");
306 }
307 Type::S32 => {
308 push_directive(tokens, "s32");
309 }
310 Type::U64 => {
311 push_directive(tokens, "u64");
312 }
313 Type::S64 => {
314 push_directive(tokens, "s64");
315 }
316 }
317 self.a.unparse_tokens(tokens);
318 tokens.push(PtxToken::Comma);
319 self.b.unparse_tokens(tokens);
320 tokens.push(PtxToken::Semicolon);
321 }
322 }
323
324}
325