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            push_opcode(tokens, "cp");
24            push_directive(tokens, "async");
25            push_directive(tokens, "ca");
26            match &self.state {
27                State::SharedCta => {
28                    push_directive(tokens, "shared::cta");
29                }
30                State::Shared => {
31                    push_directive(tokens, "shared");
32                }
33            }
34            push_directive(tokens, "global");
35            if let Some(level_cache_hint_0) = self.level_cache_hint.as_ref() {
36                match level_cache_hint_0 {
37                    LevelCacheHint::L2CacheHint => {
38                        push_directive(tokens, "L2::cache_hint");
39                    }
40                }
41            }
42            if let Some(level_prefetch_size_1) = self.level_prefetch_size.as_ref() {
43                match level_prefetch_size_1 {
44                    LevelPrefetchSize::L2128b => {
45                        push_directive(tokens, "L2::128B");
46                    }
47                    LevelPrefetchSize::L2256b => {
48                        push_directive(tokens, "L2::256B");
49                    }
50                    LevelPrefetchSize::L264b => {
51                        push_directive(tokens, "L2::64B");
52                    }
53                }
54            }
55            self.dst.unparse_tokens(tokens);
56            tokens.push(PtxToken::Comma);
57            self.src.unparse_tokens(tokens);
58            tokens.push(PtxToken::Comma);
59            match &self.cp_size {
60                CpSize::_16 => {
61                    push_token_from_str(tokens, "16");
62                }
63                CpSize::_4 => {
64                    push_token_from_str(tokens, "4");
65                }
66                CpSize::_8 => {
67                    push_token_from_str(tokens, "8");
68                }
69            }
70            if self.src_size.is_some() {
71                tokens.push(PtxToken::Comma);
72            }
73            if let Some(opt_2) = self.src_size.as_ref() {
74                opt_2.unparse_tokens(tokens);
75            }
76            if self.cache_policy.is_some() {
77                tokens.push(PtxToken::Comma);
78            }
79            if let Some(opt_3) = self.cache_policy.as_ref() {
80                opt_3.unparse_tokens(tokens);
81            }
82            tokens.push(PtxToken::Semicolon);
83        }
84    }
85
86    impl PtxUnparser for CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize {
87        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
88            push_opcode(tokens, "cp");
89            push_directive(tokens, "async");
90            push_directive(tokens, "cg");
91            match &self.state {
92                State::SharedCta => {
93                    push_directive(tokens, "shared::cta");
94                }
95                State::Shared => {
96                    push_directive(tokens, "shared");
97                }
98            }
99            push_directive(tokens, "global");
100            if let Some(level_cache_hint_4) = self.level_cache_hint.as_ref() {
101                match level_cache_hint_4 {
102                    LevelCacheHint::L2CacheHint => {
103                        push_directive(tokens, "L2::cache_hint");
104                    }
105                }
106            }
107            if let Some(level_prefetch_size_5) = self.level_prefetch_size.as_ref() {
108                match level_prefetch_size_5 {
109                    LevelPrefetchSize::L2128b => {
110                        push_directive(tokens, "L2::128B");
111                    }
112                    LevelPrefetchSize::L2256b => {
113                        push_directive(tokens, "L2::256B");
114                    }
115                    LevelPrefetchSize::L264b => {
116                        push_directive(tokens, "L2::64B");
117                    }
118                }
119            }
120            self.dst.unparse_tokens(tokens);
121            tokens.push(PtxToken::Comma);
122            self.src.unparse_tokens(tokens);
123            tokens.push(PtxToken::Comma);
124            push_token_from_str(tokens, "16");
125            if self.src_size.is_some() {
126                tokens.push(PtxToken::Comma);
127            }
128            if let Some(opt_6) = self.src_size.as_ref() {
129                opt_6.unparse_tokens(tokens);
130            }
131            if self.cache_policy.is_some() {
132                tokens.push(PtxToken::Comma);
133            }
134            if let Some(opt_7) = self.cache_policy.as_ref() {
135                opt_7.unparse_tokens(tokens);
136            }
137            tokens.push(PtxToken::Semicolon);
138        }
139    }
140
141    impl PtxUnparser for CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize1 {
142        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
143            push_opcode(tokens, "cp");
144            push_directive(tokens, "async");
145            push_directive(tokens, "ca");
146            match &self.state {
147                State::SharedCta => {
148                    push_directive(tokens, "shared::cta");
149                }
150                State::Shared => {
151                    push_directive(tokens, "shared");
152                }
153            }
154            push_directive(tokens, "global");
155            if let Some(level_cache_hint_8) = self.level_cache_hint.as_ref() {
156                match level_cache_hint_8 {
157                    LevelCacheHint::L2CacheHint => {
158                        push_directive(tokens, "L2::cache_hint");
159                    }
160                }
161            }
162            if let Some(level_prefetch_size_9) = self.level_prefetch_size.as_ref() {
163                match level_prefetch_size_9 {
164                    LevelPrefetchSize::L2128b => {
165                        push_directive(tokens, "L2::128B");
166                    }
167                    LevelPrefetchSize::L2256b => {
168                        push_directive(tokens, "L2::256B");
169                    }
170                    LevelPrefetchSize::L264b => {
171                        push_directive(tokens, "L2::64B");
172                    }
173                }
174            }
175            self.dst.unparse_tokens(tokens);
176            tokens.push(PtxToken::Comma);
177            self.src.unparse_tokens(tokens);
178            tokens.push(PtxToken::Comma);
179            match &self.cp_size {
180                CpSize::_16 => {
181                    push_token_from_str(tokens, "16");
182                }
183                CpSize::_4 => {
184                    push_token_from_str(tokens, "4");
185                }
186                CpSize::_8 => {
187                    push_token_from_str(tokens, "8");
188                }
189            }
190            if self.ignore_src.is_some() {
191                tokens.push(PtxToken::Comma);
192            }
193            if let Some(opt_10) = self.ignore_src.as_ref() {
194                opt_10.unparse_tokens(tokens);
195            }
196            if self.cache_policy.is_some() {
197                tokens.push(PtxToken::Comma);
198            }
199            if let Some(opt_11) = self.cache_policy.as_ref() {
200                opt_11.unparse_tokens(tokens);
201            }
202            tokens.push(PtxToken::Semicolon);
203        }
204    }
205
206    impl PtxUnparser for CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize1 {
207        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
208            push_opcode(tokens, "cp");
209            push_directive(tokens, "async");
210            push_directive(tokens, "cg");
211            match &self.state {
212                State::SharedCta => {
213                    push_directive(tokens, "shared::cta");
214                }
215                State::Shared => {
216                    push_directive(tokens, "shared");
217                }
218            }
219            push_directive(tokens, "global");
220            if let Some(level_cache_hint_12) = self.level_cache_hint.as_ref() {
221                match level_cache_hint_12 {
222                    LevelCacheHint::L2CacheHint => {
223                        push_directive(tokens, "L2::cache_hint");
224                    }
225                }
226            }
227            if let Some(level_prefetch_size_13) = self.level_prefetch_size.as_ref() {
228                match level_prefetch_size_13 {
229                    LevelPrefetchSize::L2128b => {
230                        push_directive(tokens, "L2::128B");
231                    }
232                    LevelPrefetchSize::L2256b => {
233                        push_directive(tokens, "L2::256B");
234                    }
235                    LevelPrefetchSize::L264b => {
236                        push_directive(tokens, "L2::64B");
237                    }
238                }
239            }
240            self.dst.unparse_tokens(tokens);
241            tokens.push(PtxToken::Comma);
242            self.src.unparse_tokens(tokens);
243            tokens.push(PtxToken::Comma);
244            push_token_from_str(tokens, "16");
245            if self.ignore_src.is_some() {
246                tokens.push(PtxToken::Comma);
247            }
248            if let Some(opt_14) = self.ignore_src.as_ref() {
249                opt_14.unparse_tokens(tokens);
250            }
251            if self.cache_policy.is_some() {
252                tokens.push(PtxToken::Comma);
253            }
254            if let Some(opt_15) = self.cache_policy.as_ref() {
255                opt_15.unparse_tokens(tokens);
256            }
257            tokens.push(PtxToken::Semicolon);
258        }
259    }
260}
ptx_parser/unparser/instruction/cp_async.rs

ptx_parser/unparser/instruction/
cp_async.rs
