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            self.unparse_tokens_mode(tokens, false);
57        }
58        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
59            push_opcode(tokens, "red");
60            push_directive(tokens, "async");
61            match &self.sem {
62                Sem::Relaxed => {
63                    push_directive(tokens, "relaxed");
64                }
65            }
66            match &self.scope {
67                Scope::Cluster => {
68                    push_directive(tokens, "cluster");
69                }
70            }
71            if let Some(ss_0) = self.ss.as_ref() {
72                match ss_0 {
73                    Ss::SharedCluster => {
74                        push_directive(tokens, "shared::cluster");
75                    }
76                }
77            }
78            match &self.completion_mechanism {
79                CompletionMechanism::MbarrierCompleteTxBytes => {
80                    push_directive(tokens, "mbarrier::complete_tx::bytes");
81                }
82            }
83            match &self.op {
84                Op::Inc => {
85                    push_directive(tokens, "inc");
86                }
87                Op::Dec => {
88                    push_directive(tokens, "dec");
89                }
90            }
91            match &self.type_ {
92                Type::U32 => {
93                    push_directive(tokens, "u32");
94                }
95            }
96            if spaced {
97                tokens.push(PtxToken::Space);
98            }
99            self.a.unparse_tokens_mode(tokens, spaced);
100            tokens.push(PtxToken::Comma);
101            if spaced {
102                tokens.push(PtxToken::Space);
103            }
104            self.b.unparse_tokens_mode(tokens, spaced);
105            tokens.push(PtxToken::Comma);
106            if spaced {
107                tokens.push(PtxToken::Space);
108            }
109            self.mbar.unparse_tokens_mode(tokens, spaced);
110            tokens.push(PtxToken::Semicolon);
111            if spaced {
112                tokens.push(PtxToken::Newline);
113            }
114        }
115    }
116}
117
118pub mod section_1 {
119    use super::*;
120    use crate::r#type::instruction::red_async::section_1::*;
121
122    impl PtxUnparser for RedAsyncSemScopeSsCompletionMechanismOpType1 {
123        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
124            self.unparse_tokens_mode(tokens, false);
125        }
126        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
127            push_opcode(tokens, "red");
128            push_directive(tokens, "async");
129            match &self.sem {
130                Sem::Relaxed => {
131                    push_directive(tokens, "relaxed");
132                }
133            }
134            match &self.scope {
135                Scope::Cluster => {
136                    push_directive(tokens, "cluster");
137                }
138            }
139            if let Some(ss_1) = self.ss.as_ref() {
140                match ss_1 {
141                    Ss::SharedCluster => {
142                        push_directive(tokens, "shared::cluster");
143                    }
144                }
145            }
146            match &self.completion_mechanism {
147                CompletionMechanism::MbarrierCompleteTxBytes => {
148                    push_directive(tokens, "mbarrier::complete_tx::bytes");
149                }
150            }
151            match &self.op {
152                Op::Min => {
153                    push_directive(tokens, "min");
154                }
155                Op::Max => {
156                    push_directive(tokens, "max");
157                }
158            }
159            match &self.type_ {
160                Type::U32 => {
161                    push_directive(tokens, "u32");
162                }
163                Type::S32 => {
164                    push_directive(tokens, "s32");
165                }
166            }
167            if spaced {
168                tokens.push(PtxToken::Space);
169            }
170            self.a.unparse_tokens_mode(tokens, spaced);
171            tokens.push(PtxToken::Comma);
172            if spaced {
173                tokens.push(PtxToken::Space);
174            }
175            self.b.unparse_tokens_mode(tokens, spaced);
176            tokens.push(PtxToken::Comma);
177            if spaced {
178                tokens.push(PtxToken::Space);
179            }
180            self.mbar.unparse_tokens_mode(tokens, spaced);
181            tokens.push(PtxToken::Semicolon);
182            if spaced {
183                tokens.push(PtxToken::Newline);
184            }
185        }
186    }
187}
188
189pub mod section_2 {
190    use super::*;
191    use crate::r#type::instruction::red_async::section_2::*;
192
193    impl PtxUnparser for RedAsyncSemScopeSsCompletionMechanismOpType2 {
194        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
195            self.unparse_tokens_mode(tokens, false);
196        }
197        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
198            push_opcode(tokens, "red");
199            push_directive(tokens, "async");
200            match &self.sem {
201                Sem::Relaxed => {
202                    push_directive(tokens, "relaxed");
203                }
204            }
205            match &self.scope {
206                Scope::Cluster => {
207                    push_directive(tokens, "cluster");
208                }
209            }
210            if let Some(ss_2) = self.ss.as_ref() {
211                match ss_2 {
212                    Ss::SharedCluster => {
213                        push_directive(tokens, "shared::cluster");
214                    }
215                }
216            }
217            match &self.completion_mechanism {
218                CompletionMechanism::MbarrierCompleteTxBytes => {
219                    push_directive(tokens, "mbarrier::complete_tx::bytes");
220                }
221            }
222            match &self.op {
223                Op::And => {
224                    push_directive(tokens, "and");
225                }
226                Op::Xor => {
227                    push_directive(tokens, "xor");
228                }
229                Op::Or => {
230                    push_directive(tokens, "or");
231                }
232            }
233            match &self.type_ {
234                Type::B32 => {
235                    push_directive(tokens, "b32");
236                }
237            }
238            if spaced {
239                tokens.push(PtxToken::Space);
240            }
241            self.a.unparse_tokens_mode(tokens, spaced);
242            tokens.push(PtxToken::Comma);
243            if spaced {
244                tokens.push(PtxToken::Space);
245            }
246            self.b.unparse_tokens_mode(tokens, spaced);
247            tokens.push(PtxToken::Comma);
248            if spaced {
249                tokens.push(PtxToken::Space);
250            }
251            self.mbar.unparse_tokens_mode(tokens, spaced);
252            tokens.push(PtxToken::Semicolon);
253            if spaced {
254                tokens.push(PtxToken::Newline);
255            }
256        }
257    }
258}
259
260pub mod section_3 {
261    use super::*;
262    use crate::r#type::instruction::red_async::section_3::*;
263
264    impl PtxUnparser for RedAsyncSemScopeSsCompletionMechanismAddType {
265        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
266            self.unparse_tokens_mode(tokens, false);
267        }
268        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
269            push_opcode(tokens, "red");
270            push_directive(tokens, "async");
271            match &self.sem {
272                Sem::Relaxed => {
273                    push_directive(tokens, "relaxed");
274                }
275            }
276            match &self.scope {
277                Scope::Cluster => {
278                    push_directive(tokens, "cluster");
279                }
280            }
281            if let Some(ss_3) = self.ss.as_ref() {
282                match ss_3 {
283                    Ss::SharedCluster => {
284                        push_directive(tokens, "shared::cluster");
285                    }
286                }
287            }
288            match &self.completion_mechanism {
289                CompletionMechanism::MbarrierCompleteTxBytes => {
290                    push_directive(tokens, "mbarrier::complete_tx::bytes");
291                }
292            }
293            push_directive(tokens, "add");
294            match &self.type_ {
295                Type::U32 => {
296                    push_directive(tokens, "u32");
297                }
298                Type::S32 => {
299                    push_directive(tokens, "s32");
300                }
301                Type::U64 => {
302                    push_directive(tokens, "u64");
303                }
304            }
305            if spaced {
306                tokens.push(PtxToken::Space);
307            }
308            self.a.unparse_tokens_mode(tokens, spaced);
309            tokens.push(PtxToken::Comma);
310            if spaced {
311                tokens.push(PtxToken::Space);
312            }
313            self.b.unparse_tokens_mode(tokens, spaced);
314            tokens.push(PtxToken::Comma);
315            if spaced {
316                tokens.push(PtxToken::Space);
317            }
318            self.mbar.unparse_tokens_mode(tokens, spaced);
319            tokens.push(PtxToken::Semicolon);
320            if spaced {
321                tokens.push(PtxToken::Newline);
322            }
323        }
324    }
325}
326
327pub mod section_4 {
328    use super::*;
329    use crate::r#type::instruction::red_async::section_4::*;
330
331    impl PtxUnparser for RedAsyncMmioSemScopeSsAddType {
332        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
333            self.unparse_tokens_mode(tokens, false);
334        }
335        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
336            push_opcode(tokens, "red");
337            push_directive(tokens, "async");
338            if self.mmio {
339                push_directive(tokens, "mmio");
340            }
341            match &self.sem {
342                Sem::Release => {
343                    push_directive(tokens, "release");
344                }
345            }
346            match &self.scope {
347                Scope::Cluster => {
348                    push_directive(tokens, "cluster");
349                }
350                Scope::Gpu => {
351                    push_directive(tokens, "gpu");
352                }
353            }
354            if let Some(ss_4) = self.ss.as_ref() {
355                match ss_4 {
356                    Ss::Global => {
357                        push_directive(tokens, "global");
358                    }
359                }
360            }
361            push_directive(tokens, "add");
362            match &self.type_ {
363                Type::U32 => {
364                    push_directive(tokens, "u32");
365                }
366                Type::S32 => {
367                    push_directive(tokens, "s32");
368                }
369                Type::U64 => {
370                    push_directive(tokens, "u64");
371                }
372                Type::S64 => {
373                    push_directive(tokens, "s64");
374                }
375            }
376            if spaced {
377                tokens.push(PtxToken::Space);
378            }
379            self.a.unparse_tokens_mode(tokens, spaced);
380            tokens.push(PtxToken::Comma);
381            if spaced {
382                tokens.push(PtxToken::Space);
383            }
384            self.b.unparse_tokens_mode(tokens, spaced);
385            tokens.push(PtxToken::Semicolon);
386            if spaced {
387                tokens.push(PtxToken::Newline);
388            }
389        }
390    }
391}
ptx_parser/unparser/instruction/red_async.rs

ptx_parser/unparser/instruction/
red_async.rs
