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

ptx_parser/unparser/instruction/
cp_async.rs
