ptx_parser/unparser/instruction/
cp_async.rs

1//! Original PTX specification:
2//!
3//! cp.async.ca.state.global{.level::cache_hint}{.level::prefetch_size} [dst], [src], cp-size{, src-size}{, cache-policy};
4//! cp.async.cg.state.global{.level::cache_hint}{.level::prefetch_size} [dst], [src], 16{, src-size}{, cache-policy};
5//! cp.async.ca.state.global{.level::cache_hint}{.level::prefetch_size} [dst], [src], cp-size{, ignore-src}{, cache-policy} ;
6//! cp.async.cg.state.global{.level::cache_hint}{.level::prefetch_size} [dst], [src], 16{, ignore-src}{, cache-policy} ;
7//! .level::cache_hint =     { .L2::cache_hint };
8//! .level::prefetch_size =  { .L2::64B, .L2::128B, .L2::256B };
9//! cp-size = { 4, 8, 16 };
10//! .state = { .shared, .shared::cta}
11
12#![allow(unused)]
13
14use crate::lexer::PtxToken;
15use crate::unparser::{PtxUnparser, common::*};
16
17pub mod section_0 {
18    use super::*;
19    use crate::r#type::instruction::cp_async::section_0::*;
20
21    impl PtxUnparser for CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize {
22        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
23            self.unparse_tokens_mode(tokens, false);
24        }
25        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
26            push_opcode(tokens, "cp");
27            push_directive(tokens, "async");
28            push_directive(tokens, "ca");
29            match &self.state {
30                State::SharedCta => {
31                    push_directive(tokens, "shared::cta");
32                }
33                State::Shared => {
34                    push_directive(tokens, "shared");
35                }
36            }
37            push_directive(tokens, "global");
38            if let Some(level_cache_hint_0) = self.level_cache_hint.as_ref() {
39                match level_cache_hint_0 {
40                    LevelCacheHint::L2CacheHint => {
41                        push_directive(tokens, "L2::cache_hint");
42                    }
43                }
44            }
45            if let Some(level_prefetch_size_1) = self.level_prefetch_size.as_ref() {
46                match level_prefetch_size_1 {
47                    LevelPrefetchSize::L2128b => {
48                        push_directive(tokens, "L2::128B");
49                    }
50                    LevelPrefetchSize::L2256b => {
51                        push_directive(tokens, "L2::256B");
52                    }
53                    LevelPrefetchSize::L264b => {
54                        push_directive(tokens, "L2::64B");
55                    }
56                }
57            }
58            if spaced {
59                tokens.push(PtxToken::Space);
60            }
61            self.dst.unparse_tokens_mode(tokens, spaced);
62            tokens.push(PtxToken::Comma);
63            if spaced {
64                tokens.push(PtxToken::Space);
65            }
66            self.src.unparse_tokens_mode(tokens, spaced);
67            tokens.push(PtxToken::Comma);
68            if spaced {
69                tokens.push(PtxToken::Space);
70            }
71            match &self.cp_size {
72                CpSize::_16 => {
73                    push_token_from_str(tokens, "16");
74                }
75                CpSize::_4 => {
76                    push_token_from_str(tokens, "4");
77                }
78                CpSize::_8 => {
79                    push_token_from_str(tokens, "8");
80                }
81            }
82            if self.src_size.is_some() {
83                tokens.push(PtxToken::Comma);
84            }
85            if let Some(opt_2) = self.src_size.as_ref() {
86                if spaced {
87                    tokens.push(PtxToken::Space);
88                }
89                opt_2.unparse_tokens_mode(tokens, spaced);
90            }
91            if self.cache_policy.is_some() {
92                tokens.push(PtxToken::Comma);
93            }
94            if let Some(opt_3) = self.cache_policy.as_ref() {
95                if spaced {
96                    tokens.push(PtxToken::Space);
97                }
98                opt_3.unparse_tokens_mode(tokens, spaced);
99            }
100            tokens.push(PtxToken::Semicolon);
101            if spaced {
102                tokens.push(PtxToken::Newline);
103            }
104        }
105    }
106
107    impl PtxUnparser for CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize {
108        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
109            self.unparse_tokens_mode(tokens, false);
110        }
111        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
112            push_opcode(tokens, "cp");
113            push_directive(tokens, "async");
114            push_directive(tokens, "cg");
115            match &self.state {
116                State::SharedCta => {
117                    push_directive(tokens, "shared::cta");
118                }
119                State::Shared => {
120                    push_directive(tokens, "shared");
121                }
122            }
123            push_directive(tokens, "global");
124            if let Some(level_cache_hint_4) = self.level_cache_hint.as_ref() {
125                match level_cache_hint_4 {
126                    LevelCacheHint::L2CacheHint => {
127                        push_directive(tokens, "L2::cache_hint");
128                    }
129                }
130            }
131            if let Some(level_prefetch_size_5) = self.level_prefetch_size.as_ref() {
132                match level_prefetch_size_5 {
133                    LevelPrefetchSize::L2128b => {
134                        push_directive(tokens, "L2::128B");
135                    }
136                    LevelPrefetchSize::L2256b => {
137                        push_directive(tokens, "L2::256B");
138                    }
139                    LevelPrefetchSize::L264b => {
140                        push_directive(tokens, "L2::64B");
141                    }
142                }
143            }
144            if spaced {
145                tokens.push(PtxToken::Space);
146            }
147            self.dst.unparse_tokens_mode(tokens, spaced);
148            tokens.push(PtxToken::Comma);
149            if spaced {
150                tokens.push(PtxToken::Space);
151            }
152            self.src.unparse_tokens_mode(tokens, spaced);
153            tokens.push(PtxToken::Comma);
154            if spaced {
155                tokens.push(PtxToken::Space);
156            }
157            push_token_from_str(tokens, "16");
158            if self.src_size.is_some() {
159                tokens.push(PtxToken::Comma);
160            }
161            if let Some(opt_6) = self.src_size.as_ref() {
162                if spaced {
163                    tokens.push(PtxToken::Space);
164                }
165                opt_6.unparse_tokens_mode(tokens, spaced);
166            }
167            if self.cache_policy.is_some() {
168                tokens.push(PtxToken::Comma);
169            }
170            if let Some(opt_7) = self.cache_policy.as_ref() {
171                if spaced {
172                    tokens.push(PtxToken::Space);
173                }
174                opt_7.unparse_tokens_mode(tokens, spaced);
175            }
176            tokens.push(PtxToken::Semicolon);
177            if spaced {
178                tokens.push(PtxToken::Newline);
179            }
180        }
181    }
182
183    impl PtxUnparser for CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize1 {
184        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
185            self.unparse_tokens_mode(tokens, false);
186        }
187        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
188            push_opcode(tokens, "cp");
189            push_directive(tokens, "async");
190            push_directive(tokens, "ca");
191            match &self.state {
192                State::SharedCta => {
193                    push_directive(tokens, "shared::cta");
194                }
195                State::Shared => {
196                    push_directive(tokens, "shared");
197                }
198            }
199            push_directive(tokens, "global");
200            if let Some(level_cache_hint_8) = self.level_cache_hint.as_ref() {
201                match level_cache_hint_8 {
202                    LevelCacheHint::L2CacheHint => {
203                        push_directive(tokens, "L2::cache_hint");
204                    }
205                }
206            }
207            if let Some(level_prefetch_size_9) = self.level_prefetch_size.as_ref() {
208                match level_prefetch_size_9 {
209                    LevelPrefetchSize::L2128b => {
210                        push_directive(tokens, "L2::128B");
211                    }
212                    LevelPrefetchSize::L2256b => {
213                        push_directive(tokens, "L2::256B");
214                    }
215                    LevelPrefetchSize::L264b => {
216                        push_directive(tokens, "L2::64B");
217                    }
218                }
219            }
220            if spaced {
221                tokens.push(PtxToken::Space);
222            }
223            self.dst.unparse_tokens_mode(tokens, spaced);
224            tokens.push(PtxToken::Comma);
225            if spaced {
226                tokens.push(PtxToken::Space);
227            }
228            self.src.unparse_tokens_mode(tokens, spaced);
229            tokens.push(PtxToken::Comma);
230            if spaced {
231                tokens.push(PtxToken::Space);
232            }
233            match &self.cp_size {
234                CpSize::_16 => {
235                    push_token_from_str(tokens, "16");
236                }
237                CpSize::_4 => {
238                    push_token_from_str(tokens, "4");
239                }
240                CpSize::_8 => {
241                    push_token_from_str(tokens, "8");
242                }
243            }
244            if self.ignore_src.is_some() {
245                tokens.push(PtxToken::Comma);
246            }
247            if let Some(opt_10) = self.ignore_src.as_ref() {
248                if spaced {
249                    tokens.push(PtxToken::Space);
250                }
251                opt_10.unparse_tokens_mode(tokens, spaced);
252            }
253            if self.cache_policy.is_some() {
254                tokens.push(PtxToken::Comma);
255            }
256            if let Some(opt_11) = self.cache_policy.as_ref() {
257                if spaced {
258                    tokens.push(PtxToken::Space);
259                }
260                opt_11.unparse_tokens_mode(tokens, spaced);
261            }
262            tokens.push(PtxToken::Semicolon);
263            if spaced {
264                tokens.push(PtxToken::Newline);
265            }
266        }
267    }
268
269    impl PtxUnparser for CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize1 {
270        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
271            self.unparse_tokens_mode(tokens, false);
272        }
273        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
274            push_opcode(tokens, "cp");
275            push_directive(tokens, "async");
276            push_directive(tokens, "cg");
277            match &self.state {
278                State::SharedCta => {
279                    push_directive(tokens, "shared::cta");
280                }
281                State::Shared => {
282                    push_directive(tokens, "shared");
283                }
284            }
285            push_directive(tokens, "global");
286            if let Some(level_cache_hint_12) = self.level_cache_hint.as_ref() {
287                match level_cache_hint_12 {
288                    LevelCacheHint::L2CacheHint => {
289                        push_directive(tokens, "L2::cache_hint");
290                    }
291                }
292            }
293            if let Some(level_prefetch_size_13) = self.level_prefetch_size.as_ref() {
294                match level_prefetch_size_13 {
295                    LevelPrefetchSize::L2128b => {
296                        push_directive(tokens, "L2::128B");
297                    }
298                    LevelPrefetchSize::L2256b => {
299                        push_directive(tokens, "L2::256B");
300                    }
301                    LevelPrefetchSize::L264b => {
302                        push_directive(tokens, "L2::64B");
303                    }
304                }
305            }
306            if spaced {
307                tokens.push(PtxToken::Space);
308            }
309            self.dst.unparse_tokens_mode(tokens, spaced);
310            tokens.push(PtxToken::Comma);
311            if spaced {
312                tokens.push(PtxToken::Space);
313            }
314            self.src.unparse_tokens_mode(tokens, spaced);
315            tokens.push(PtxToken::Comma);
316            if spaced {
317                tokens.push(PtxToken::Space);
318            }
319            push_token_from_str(tokens, "16");
320            if self.ignore_src.is_some() {
321                tokens.push(PtxToken::Comma);
322            }
323            if let Some(opt_14) = self.ignore_src.as_ref() {
324                if spaced {
325                    tokens.push(PtxToken::Space);
326                }
327                opt_14.unparse_tokens_mode(tokens, spaced);
328            }
329            if self.cache_policy.is_some() {
330                tokens.push(PtxToken::Comma);
331            }
332            if let Some(opt_15) = self.cache_policy.as_ref() {
333                if spaced {
334                    tokens.push(PtxToken::Space);
335                }
336                opt_15.unparse_tokens_mode(tokens, spaced);
337            }
338            tokens.push(PtxToken::Semicolon);
339            if spaced {
340                tokens.push(PtxToken::Newline);
341            }
342        }
343    }
344}
ptx_parser/unparser/instruction/cp_async.rs

ptx_parser/unparser/instruction/
cp_async.rs
