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
103pub mod section_1 {
104    use super::*;
105    use crate::r#type::instruction::red_async::section_1::*;
106
107    impl PtxUnparser for RedAsyncSemScopeSsCompletionMechanismOpType1 {
108        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
109            push_opcode(tokens, "red");
110            push_directive(tokens, "async");
111            match &self.sem {
112                Sem::Relaxed => {
113                    push_directive(tokens, "relaxed");
114                }
115            }
116            match &self.scope {
117                Scope::Cluster => {
118                    push_directive(tokens, "cluster");
119                }
120            }
121            if let Some(ss_1) = self.ss.as_ref() {
122                match ss_1 {
123                    Ss::SharedCluster => {
124                        push_directive(tokens, "shared::cluster");
125                    }
126                }
127            }
128            match &self.completion_mechanism {
129                CompletionMechanism::MbarrierCompleteTxBytes => {
130                    push_directive(tokens, "mbarrier::complete_tx::bytes");
131                }
132            }
133            match &self.op {
134                Op::Min => {
135                    push_directive(tokens, "min");
136                }
137                Op::Max => {
138                    push_directive(tokens, "max");
139                }
140            }
141            match &self.type_ {
142                Type::U32 => {
143                    push_directive(tokens, "u32");
144                }
145                Type::S32 => {
146                    push_directive(tokens, "s32");
147                }
148            }
149            self.a.unparse_tokens(tokens);
150            tokens.push(PtxToken::Comma);
151            self.b.unparse_tokens(tokens);
152            tokens.push(PtxToken::Comma);
153            self.mbar.unparse_tokens(tokens);
154            tokens.push(PtxToken::Semicolon);
155        }
156    }
157}
158
159pub mod section_2 {
160    use super::*;
161    use crate::r#type::instruction::red_async::section_2::*;
162
163    impl PtxUnparser for RedAsyncSemScopeSsCompletionMechanismOpType2 {
164        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
165            push_opcode(tokens, "red");
166            push_directive(tokens, "async");
167            match &self.sem {
168                Sem::Relaxed => {
169                    push_directive(tokens, "relaxed");
170                }
171            }
172            match &self.scope {
173                Scope::Cluster => {
174                    push_directive(tokens, "cluster");
175                }
176            }
177            if let Some(ss_2) = self.ss.as_ref() {
178                match ss_2 {
179                    Ss::SharedCluster => {
180                        push_directive(tokens, "shared::cluster");
181                    }
182                }
183            }
184            match &self.completion_mechanism {
185                CompletionMechanism::MbarrierCompleteTxBytes => {
186                    push_directive(tokens, "mbarrier::complete_tx::bytes");
187                }
188            }
189            match &self.op {
190                Op::And => {
191                    push_directive(tokens, "and");
192                }
193                Op::Xor => {
194                    push_directive(tokens, "xor");
195                }
196                Op::Or => {
197                    push_directive(tokens, "or");
198                }
199            }
200            match &self.type_ {
201                Type::B32 => {
202                    push_directive(tokens, "b32");
203                }
204            }
205            self.a.unparse_tokens(tokens);
206            tokens.push(PtxToken::Comma);
207            self.b.unparse_tokens(tokens);
208            tokens.push(PtxToken::Comma);
209            self.mbar.unparse_tokens(tokens);
210            tokens.push(PtxToken::Semicolon);
211        }
212    }
213}
214
215pub mod section_3 {
216    use super::*;
217    use crate::r#type::instruction::red_async::section_3::*;
218
219    impl PtxUnparser for RedAsyncSemScopeSsCompletionMechanismAddType {
220        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
221            push_opcode(tokens, "red");
222            push_directive(tokens, "async");
223            match &self.sem {
224                Sem::Relaxed => {
225                    push_directive(tokens, "relaxed");
226                }
227            }
228            match &self.scope {
229                Scope::Cluster => {
230                    push_directive(tokens, "cluster");
231                }
232            }
233            if let Some(ss_3) = self.ss.as_ref() {
234                match ss_3 {
235                    Ss::SharedCluster => {
236                        push_directive(tokens, "shared::cluster");
237                    }
238                }
239            }
240            match &self.completion_mechanism {
241                CompletionMechanism::MbarrierCompleteTxBytes => {
242                    push_directive(tokens, "mbarrier::complete_tx::bytes");
243                }
244            }
245            push_directive(tokens, "add");
246            match &self.type_ {
247                Type::U32 => {
248                    push_directive(tokens, "u32");
249                }
250                Type::S32 => {
251                    push_directive(tokens, "s32");
252                }
253                Type::U64 => {
254                    push_directive(tokens, "u64");
255                }
256            }
257            self.a.unparse_tokens(tokens);
258            tokens.push(PtxToken::Comma);
259            self.b.unparse_tokens(tokens);
260            tokens.push(PtxToken::Comma);
261            self.mbar.unparse_tokens(tokens);
262            tokens.push(PtxToken::Semicolon);
263        }
264    }
265}
266
267pub mod section_4 {
268    use super::*;
269    use crate::r#type::instruction::red_async::section_4::*;
270
271    impl PtxUnparser for RedAsyncMmioSemScopeSsAddType {
272        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
273            push_opcode(tokens, "red");
274            push_directive(tokens, "async");
275            if self.mmio {
276                push_directive(tokens, "mmio");
277            }
278            match &self.sem {
279                Sem::Release => {
280                    push_directive(tokens, "release");
281                }
282            }
283            match &self.scope {
284                Scope::Cluster => {
285                    push_directive(tokens, "cluster");
286                }
287                Scope::Gpu => {
288                    push_directive(tokens, "gpu");
289                }
290            }
291            if let Some(ss_4) = self.ss.as_ref() {
292                match ss_4 {
293                    Ss::Global => {
294                        push_directive(tokens, "global");
295                    }
296                }
297            }
298            push_directive(tokens, "add");
299            match &self.type_ {
300                Type::U32 => {
301                    push_directive(tokens, "u32");
302                }
303                Type::S32 => {
304                    push_directive(tokens, "s32");
305                }
306                Type::U64 => {
307                    push_directive(tokens, "u64");
308                }
309                Type::S64 => {
310                    push_directive(tokens, "s64");
311                }
312            }
313            self.a.unparse_tokens(tokens);
314            tokens.push(PtxToken::Comma);
315            self.b.unparse_tokens(tokens);
316            tokens.push(PtxToken::Semicolon);
317        }
318    }
319}
ptx_parser/unparser/instruction/red_async.rs

ptx_parser/unparser/instruction/
red_async.rs
