ptx_parser/unparser/instruction/
cp_async_bulk.rs

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

ptx_parser/unparser/instruction/
cp_async_bulk.rs
