ptx_parser/unparser/instruction/
cp_reduce_async_bulk.rs

1//! Original PTX specification:
2//!
3//! cp.reduce.async.bulk.dst.src.completion_mechanism.redOp.type [dstMem], [srcMem], size, [mbar];
4//! .dst =                  { .shared::cluster };
5//! .src =                  { .shared::cta };
6//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
7//! .redOp=                 { .and, .or, .xor, .add, .inc, .dec, .min, .max };
8//! .type =                 { .b32, .u32, .s32, .b64, .u64 };
9//! ----------------------------------------------------------------
10//! cp.reduce.async.bulk.dst.src.completion_mechanism{.level::cache_hint}.redOp.type [dstMem], [srcMem], size{, cache-policy};
11//! .dst =                  { .global      };
12//! .src =                  { .shared::cta };
13//! ----------------------------------------------------------------
14//! .completion_mechanism = { .bulk_group };
15//! .level::cache_hint    = { .L2::cache_hint };
16//! .redOp=                 { .and, .or, .xor, .add, .inc, .dec, .min, .max };
17//! .type =                 { .f16, .bf16, .b32, .u32, .s32, .b64, .u64, .s64, .f32, .f64 };
18//! ----------------------------------------------------------------
19//! cp.reduce.async.bulk.dst.src.completion_mechanism{.level::cache_hint}.add.noftz.type [dstMem], [srcMem], size{, cache-policy};
20//! .dst  =                 { .global };
21//! .src  =                 { .shared::cta };
22//! .completion_mechanism = { .bulk_group };
23//! .type =                 { .f16, .bf16 };
24
25#![allow(unused)]
26
27use crate::lexer::PtxToken;
28use crate::unparser::{PtxUnparser, common::*};
29
30pub mod section_0 {
31    use super::*;
32    use crate::r#type::instruction::cp_reduce_async_bulk::section_0::*;
33
34    impl PtxUnparser for CpReduceAsyncBulkDstSrcCompletionMechanismRedopType {
35        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
36            self.unparse_tokens_mode(tokens, false);
37        }
38        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
39            push_opcode(tokens, "cp");
40            push_directive(tokens, "reduce");
41            push_directive(tokens, "async");
42            push_directive(tokens, "bulk");
43            match &self.dst {
44                Dst::SharedCluster => {
45                    push_directive(tokens, "shared::cluster");
46                }
47            }
48            match &self.src {
49                Src::SharedCta => {
50                    push_directive(tokens, "shared::cta");
51                }
52            }
53            match &self.completion_mechanism {
54                CompletionMechanism::MbarrierCompleteTxBytes => {
55                    push_directive(tokens, "mbarrier::complete_tx::bytes");
56                }
57            }
58            match &self.redop {
59                Redop::And => {
60                    push_directive(tokens, "and");
61                }
62                Redop::Xor => {
63                    push_directive(tokens, "xor");
64                }
65                Redop::Add => {
66                    push_directive(tokens, "add");
67                }
68                Redop::Inc => {
69                    push_directive(tokens, "inc");
70                }
71                Redop::Dec => {
72                    push_directive(tokens, "dec");
73                }
74                Redop::Min => {
75                    push_directive(tokens, "min");
76                }
77                Redop::Max => {
78                    push_directive(tokens, "max");
79                }
80                Redop::Or => {
81                    push_directive(tokens, "or");
82                }
83            }
84            match &self.type_ {
85                Type::B32 => {
86                    push_directive(tokens, "b32");
87                }
88                Type::U32 => {
89                    push_directive(tokens, "u32");
90                }
91                Type::S32 => {
92                    push_directive(tokens, "s32");
93                }
94                Type::B64 => {
95                    push_directive(tokens, "b64");
96                }
97                Type::U64 => {
98                    push_directive(tokens, "u64");
99                }
100            }
101            if spaced {
102                tokens.push(PtxToken::Space);
103            }
104            self.dstmem.unparse_tokens_mode(tokens, spaced);
105            tokens.push(PtxToken::Comma);
106            if spaced {
107                tokens.push(PtxToken::Space);
108            }
109            self.srcmem.unparse_tokens_mode(tokens, spaced);
110            tokens.push(PtxToken::Comma);
111            if spaced {
112                tokens.push(PtxToken::Space);
113            }
114            self.size.unparse_tokens_mode(tokens, spaced);
115            tokens.push(PtxToken::Comma);
116            if spaced {
117                tokens.push(PtxToken::Space);
118            }
119            self.mbar.unparse_tokens_mode(tokens, spaced);
120            tokens.push(PtxToken::Semicolon);
121            if spaced {
122                tokens.push(PtxToken::Newline);
123            }
124        }
125    }
126}
127
128pub mod section_1 {
129    use super::*;
130    use crate::r#type::instruction::cp_reduce_async_bulk::section_1::*;
131
132    impl PtxUnparser for CpReduceAsyncBulkDstSrcCompletionMechanismLevelCacheHintRedopType {
133        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
134            self.unparse_tokens_mode(tokens, false);
135        }
136        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
137            push_opcode(tokens, "cp");
138            push_directive(tokens, "reduce");
139            push_directive(tokens, "async");
140            push_directive(tokens, "bulk");
141            match &self.dst {
142                Dst::Global => {
143                    push_directive(tokens, "global");
144                }
145            }
146            match &self.src {
147                Src::SharedCta => {
148                    push_directive(tokens, "shared::cta");
149                }
150            }
151            match &self.completion_mechanism {
152                CompletionMechanism::MbarrierCompleteTxBytes => {
153                    push_directive(tokens, "mbarrier::complete_tx::bytes");
154                }
155            }
156            if self.level_cache_hint {
157                push_directive(tokens, "level::cache_hint");
158            }
159            match &self.redop {
160                Redop::And => {
161                    push_directive(tokens, "and");
162                }
163                Redop::Xor => {
164                    push_directive(tokens, "xor");
165                }
166                Redop::Add => {
167                    push_directive(tokens, "add");
168                }
169                Redop::Inc => {
170                    push_directive(tokens, "inc");
171                }
172                Redop::Dec => {
173                    push_directive(tokens, "dec");
174                }
175                Redop::Min => {
176                    push_directive(tokens, "min");
177                }
178                Redop::Max => {
179                    push_directive(tokens, "max");
180                }
181                Redop::Or => {
182                    push_directive(tokens, "or");
183                }
184            }
185            match &self.type_ {
186                Type::B32 => {
187                    push_directive(tokens, "b32");
188                }
189                Type::U32 => {
190                    push_directive(tokens, "u32");
191                }
192                Type::S32 => {
193                    push_directive(tokens, "s32");
194                }
195                Type::B64 => {
196                    push_directive(tokens, "b64");
197                }
198                Type::U64 => {
199                    push_directive(tokens, "u64");
200                }
201            }
202            if spaced {
203                tokens.push(PtxToken::Space);
204            }
205            self.dstmem.unparse_tokens_mode(tokens, spaced);
206            tokens.push(PtxToken::Comma);
207            if spaced {
208                tokens.push(PtxToken::Space);
209            }
210            self.srcmem.unparse_tokens_mode(tokens, spaced);
211            tokens.push(PtxToken::Comma);
212            if spaced {
213                tokens.push(PtxToken::Space);
214            }
215            self.size.unparse_tokens_mode(tokens, spaced);
216            if self.cache_policy.is_some() {
217                tokens.push(PtxToken::Comma);
218            }
219            if let Some(opt_0) = self.cache_policy.as_ref() {
220                if spaced {
221                    tokens.push(PtxToken::Space);
222                }
223                opt_0.unparse_tokens_mode(tokens, spaced);
224            }
225            tokens.push(PtxToken::Semicolon);
226            if spaced {
227                tokens.push(PtxToken::Newline);
228            }
229        }
230    }
231}
232
233pub mod section_2 {
234    use super::*;
235    use crate::r#type::instruction::cp_reduce_async_bulk::section_2::*;
236
237    impl PtxUnparser for CpReduceAsyncBulkDstSrcCompletionMechanismLevelCacheHintAddNoftzType {
238        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
239            self.unparse_tokens_mode(tokens, false);
240        }
241        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
242            push_opcode(tokens, "cp");
243            push_directive(tokens, "reduce");
244            push_directive(tokens, "async");
245            push_directive(tokens, "bulk");
246            match &self.dst {
247                Dst::Global => {
248                    push_directive(tokens, "global");
249                }
250            }
251            match &self.src {
252                Src::SharedCta => {
253                    push_directive(tokens, "shared::cta");
254                }
255            }
256            match &self.completion_mechanism {
257                CompletionMechanism::BulkGroup => {
258                    push_directive(tokens, "bulk_group");
259                }
260            }
261            if let Some(level_cache_hint_1) = self.level_cache_hint.as_ref() {
262                match level_cache_hint_1 {
263                    LevelCacheHint::L2CacheHint => {
264                        push_directive(tokens, "L2::cache_hint");
265                    }
266                }
267            }
268            push_directive(tokens, "add");
269            push_directive(tokens, "noftz");
270            match &self.type_ {
271                Type::Bf16 => {
272                    push_directive(tokens, "bf16");
273                }
274                Type::F16 => {
275                    push_directive(tokens, "f16");
276                }
277            }
278            if spaced {
279                tokens.push(PtxToken::Space);
280            }
281            self.dstmem.unparse_tokens_mode(tokens, spaced);
282            tokens.push(PtxToken::Comma);
283            if spaced {
284                tokens.push(PtxToken::Space);
285            }
286            self.srcmem.unparse_tokens_mode(tokens, spaced);
287            tokens.push(PtxToken::Comma);
288            if spaced {
289                tokens.push(PtxToken::Space);
290            }
291            self.size.unparse_tokens_mode(tokens, spaced);
292            if self.cache_policy.is_some() {
293                tokens.push(PtxToken::Comma);
294            }
295            if let Some(opt_2) = self.cache_policy.as_ref() {
296                if spaced {
297                    tokens.push(PtxToken::Space);
298                }
299                opt_2.unparse_tokens_mode(tokens, spaced);
300            }
301            tokens.push(PtxToken::Semicolon);
302            if spaced {
303                tokens.push(PtxToken::Newline);
304            }
305        }
306    }
307}
ptx_parser/unparser/instruction/cp_reduce_async_bulk.rs

ptx_parser/unparser/instruction/
cp_reduce_async_bulk.rs
