ptx_parser/unparser/instruction/ld_global_nc.rs
1//! Original PTX specification:
2//!
3//! ld.global{.cop}.nc{.level::cache_hint}{.level::prefetch_size}.type d, [a]{, cache-policy};
4//! ld.global{.cop}.nc{.level::cache_hint}{.level::prefetch_size}.vec.type d, [a]{, cache-policy};
5//! ld.global.nc{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.level::prefetch_size}.type d, [a]{, cache-policy};
6//! ld.global.nc{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.level::prefetch_size}.vec.type d, [a]{, cache-policy};
7//! .cop = { .ca, .cg, .cs }; // cache operation
8//! .level1::eviction_priority = { .L1::evict_normal, .L1::evict_unchanged, .L1::evict_first, .L1::evict_last, .L1::no_allocate};
9//! .level2::eviction_priority = {.L2::evict_normal, .L2::evict_first, .L2::evict_last};
10//! .level::cache_hint = { .L2::cache_hint };
11//! .level::prefetch_size = { .L2::64B, .L2::128B, .L2::256B };
12//! .vec = { .v2, .v4, .v8 };
13//! .type = { .b8, .b16, .b32, .b64, .b128,
14//! .u8, .u16, .u32, .u64,
15//! .s8, .s16, .s32, .s64,
16//! .f32, .f64 };
17
18#![allow(unused)]
19
20use crate::lexer::PtxToken;
21use crate::unparser::{PtxUnparser, common::*};
22
23pub mod section_0 {
24 use super::*;
25 use crate::r#type::instruction::ld_global_nc::section_0::*;
26
27 impl PtxUnparser for LdGlobalCopNcLevelCacheHintLevelPrefetchSizeType {
28 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
29 push_opcode(tokens, "ld");
30 push_directive(tokens, "global");
31 if let Some(cop_0) = self.cop.as_ref() {
32 match cop_0 {
33 Cop::Ca => {
34 push_directive(tokens, "ca");
35 }
36 Cop::Cg => {
37 push_directive(tokens, "cg");
38 }
39 Cop::Cs => {
40 push_directive(tokens, "cs");
41 }
42 }
43 }
44 push_directive(tokens, "nc");
45 if let Some(level_cache_hint_1) = self.level_cache_hint.as_ref() {
46 match level_cache_hint_1 {
47 LevelCacheHint::L2CacheHint => {
48 push_directive(tokens, "L2::cache_hint");
49 }
50 }
51 }
52 if let Some(level_prefetch_size_2) = self.level_prefetch_size.as_ref() {
53 match level_prefetch_size_2 {
54 LevelPrefetchSize::L2128b => {
55 push_directive(tokens, "L2::128B");
56 }
57 LevelPrefetchSize::L2256b => {
58 push_directive(tokens, "L2::256B");
59 }
60 LevelPrefetchSize::L264b => {
61 push_directive(tokens, "L2::64B");
62 }
63 }
64 }
65 match &self.type_ {
66 Type::B128 => {
67 push_directive(tokens, "b128");
68 }
69 Type::B16 => {
70 push_directive(tokens, "b16");
71 }
72 Type::B32 => {
73 push_directive(tokens, "b32");
74 }
75 Type::B64 => {
76 push_directive(tokens, "b64");
77 }
78 Type::U16 => {
79 push_directive(tokens, "u16");
80 }
81 Type::U32 => {
82 push_directive(tokens, "u32");
83 }
84 Type::U64 => {
85 push_directive(tokens, "u64");
86 }
87 Type::S16 => {
88 push_directive(tokens, "s16");
89 }
90 Type::S32 => {
91 push_directive(tokens, "s32");
92 }
93 Type::S64 => {
94 push_directive(tokens, "s64");
95 }
96 Type::F32 => {
97 push_directive(tokens, "f32");
98 }
99 Type::F64 => {
100 push_directive(tokens, "f64");
101 }
102 Type::B8 => {
103 push_directive(tokens, "b8");
104 }
105 Type::U8 => {
106 push_directive(tokens, "u8");
107 }
108 Type::S8 => {
109 push_directive(tokens, "s8");
110 }
111 }
112 self.d.unparse_tokens(tokens);
113 tokens.push(PtxToken::Comma);
114 self.a.unparse_tokens(tokens);
115 if self.cache_policy.is_some() { tokens.push(PtxToken::Comma); }
116 if let Some(opt_3) = self.cache_policy.as_ref() {
117 opt_3.unparse_tokens(tokens);
118 }
119 tokens.push(PtxToken::Semicolon);
120 }
121 }
122
123 impl PtxUnparser for LdGlobalCopNcLevelCacheHintLevelPrefetchSizeVecType {
124 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
125 push_opcode(tokens, "ld");
126 push_directive(tokens, "global");
127 if let Some(cop_4) = self.cop.as_ref() {
128 match cop_4 {
129 Cop::Ca => {
130 push_directive(tokens, "ca");
131 }
132 Cop::Cg => {
133 push_directive(tokens, "cg");
134 }
135 Cop::Cs => {
136 push_directive(tokens, "cs");
137 }
138 }
139 }
140 push_directive(tokens, "nc");
141 if let Some(level_cache_hint_5) = self.level_cache_hint.as_ref() {
142 match level_cache_hint_5 {
143 LevelCacheHint::L2CacheHint => {
144 push_directive(tokens, "L2::cache_hint");
145 }
146 }
147 }
148 if let Some(level_prefetch_size_6) = self.level_prefetch_size.as_ref() {
149 match level_prefetch_size_6 {
150 LevelPrefetchSize::L2128b => {
151 push_directive(tokens, "L2::128B");
152 }
153 LevelPrefetchSize::L2256b => {
154 push_directive(tokens, "L2::256B");
155 }
156 LevelPrefetchSize::L264b => {
157 push_directive(tokens, "L2::64B");
158 }
159 }
160 }
161 match &self.vec {
162 Vec::V2 => {
163 push_directive(tokens, "v2");
164 }
165 Vec::V4 => {
166 push_directive(tokens, "v4");
167 }
168 Vec::V8 => {
169 push_directive(tokens, "v8");
170 }
171 }
172 match &self.type_ {
173 Type::B128 => {
174 push_directive(tokens, "b128");
175 }
176 Type::B16 => {
177 push_directive(tokens, "b16");
178 }
179 Type::B32 => {
180 push_directive(tokens, "b32");
181 }
182 Type::B64 => {
183 push_directive(tokens, "b64");
184 }
185 Type::U16 => {
186 push_directive(tokens, "u16");
187 }
188 Type::U32 => {
189 push_directive(tokens, "u32");
190 }
191 Type::U64 => {
192 push_directive(tokens, "u64");
193 }
194 Type::S16 => {
195 push_directive(tokens, "s16");
196 }
197 Type::S32 => {
198 push_directive(tokens, "s32");
199 }
200 Type::S64 => {
201 push_directive(tokens, "s64");
202 }
203 Type::F32 => {
204 push_directive(tokens, "f32");
205 }
206 Type::F64 => {
207 push_directive(tokens, "f64");
208 }
209 Type::B8 => {
210 push_directive(tokens, "b8");
211 }
212 Type::U8 => {
213 push_directive(tokens, "u8");
214 }
215 Type::S8 => {
216 push_directive(tokens, "s8");
217 }
218 }
219 self.d.unparse_tokens(tokens);
220 tokens.push(PtxToken::Comma);
221 self.a.unparse_tokens(tokens);
222 if self.cache_policy.is_some() { tokens.push(PtxToken::Comma); }
223 if let Some(opt_7) = self.cache_policy.as_ref() {
224 opt_7.unparse_tokens(tokens);
225 }
226 tokens.push(PtxToken::Semicolon);
227 }
228 }
229
230 impl PtxUnparser for LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeType {
231 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
232 push_opcode(tokens, "ld");
233 push_directive(tokens, "global");
234 push_directive(tokens, "nc");
235 if let Some(level1_eviction_priority_8) = self.level1_eviction_priority.as_ref() {
236 match level1_eviction_priority_8 {
237 Level1EvictionPriority::L1EvictUnchanged => {
238 push_directive(tokens, "L1::evict_unchanged");
239 }
240 Level1EvictionPriority::L1EvictNormal => {
241 push_directive(tokens, "L1::evict_normal");
242 }
243 Level1EvictionPriority::L1EvictFirst => {
244 push_directive(tokens, "L1::evict_first");
245 }
246 Level1EvictionPriority::L1NoAllocate => {
247 push_directive(tokens, "L1::no_allocate");
248 }
249 Level1EvictionPriority::L1EvictLast => {
250 push_directive(tokens, "L1::evict_last");
251 }
252 }
253 }
254 if let Some(level2_eviction_priority_9) = self.level2_eviction_priority.as_ref() {
255 match level2_eviction_priority_9 {
256 Level2EvictionPriority::L2EvictNormal => {
257 push_directive(tokens, "L2::evict_normal");
258 }
259 Level2EvictionPriority::L2EvictFirst => {
260 push_directive(tokens, "L2::evict_first");
261 }
262 Level2EvictionPriority::L2EvictLast => {
263 push_directive(tokens, "L2::evict_last");
264 }
265 }
266 }
267 if let Some(level_cache_hint_10) = self.level_cache_hint.as_ref() {
268 match level_cache_hint_10 {
269 LevelCacheHint::L2CacheHint => {
270 push_directive(tokens, "L2::cache_hint");
271 }
272 }
273 }
274 if let Some(level_prefetch_size_11) = self.level_prefetch_size.as_ref() {
275 match level_prefetch_size_11 {
276 LevelPrefetchSize::L2128b => {
277 push_directive(tokens, "L2::128B");
278 }
279 LevelPrefetchSize::L2256b => {
280 push_directive(tokens, "L2::256B");
281 }
282 LevelPrefetchSize::L264b => {
283 push_directive(tokens, "L2::64B");
284 }
285 }
286 }
287 match &self.type_ {
288 Type::B128 => {
289 push_directive(tokens, "b128");
290 }
291 Type::B16 => {
292 push_directive(tokens, "b16");
293 }
294 Type::B32 => {
295 push_directive(tokens, "b32");
296 }
297 Type::B64 => {
298 push_directive(tokens, "b64");
299 }
300 Type::U16 => {
301 push_directive(tokens, "u16");
302 }
303 Type::U32 => {
304 push_directive(tokens, "u32");
305 }
306 Type::U64 => {
307 push_directive(tokens, "u64");
308 }
309 Type::S16 => {
310 push_directive(tokens, "s16");
311 }
312 Type::S32 => {
313 push_directive(tokens, "s32");
314 }
315 Type::S64 => {
316 push_directive(tokens, "s64");
317 }
318 Type::F32 => {
319 push_directive(tokens, "f32");
320 }
321 Type::F64 => {
322 push_directive(tokens, "f64");
323 }
324 Type::B8 => {
325 push_directive(tokens, "b8");
326 }
327 Type::U8 => {
328 push_directive(tokens, "u8");
329 }
330 Type::S8 => {
331 push_directive(tokens, "s8");
332 }
333 }
334 self.d.unparse_tokens(tokens);
335 tokens.push(PtxToken::Comma);
336 self.a.unparse_tokens(tokens);
337 if self.cache_policy.is_some() { tokens.push(PtxToken::Comma); }
338 if let Some(opt_12) = self.cache_policy.as_ref() {
339 opt_12.unparse_tokens(tokens);
340 }
341 tokens.push(PtxToken::Semicolon);
342 }
343 }
344
345 impl PtxUnparser for LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType {
346 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
347 push_opcode(tokens, "ld");
348 push_directive(tokens, "global");
349 push_directive(tokens, "nc");
350 if let Some(level1_eviction_priority_13) = self.level1_eviction_priority.as_ref() {
351 match level1_eviction_priority_13 {
352 Level1EvictionPriority::L1EvictUnchanged => {
353 push_directive(tokens, "L1::evict_unchanged");
354 }
355 Level1EvictionPriority::L1EvictNormal => {
356 push_directive(tokens, "L1::evict_normal");
357 }
358 Level1EvictionPriority::L1EvictFirst => {
359 push_directive(tokens, "L1::evict_first");
360 }
361 Level1EvictionPriority::L1NoAllocate => {
362 push_directive(tokens, "L1::no_allocate");
363 }
364 Level1EvictionPriority::L1EvictLast => {
365 push_directive(tokens, "L1::evict_last");
366 }
367 }
368 }
369 if let Some(level2_eviction_priority_14) = self.level2_eviction_priority.as_ref() {
370 match level2_eviction_priority_14 {
371 Level2EvictionPriority::L2EvictNormal => {
372 push_directive(tokens, "L2::evict_normal");
373 }
374 Level2EvictionPriority::L2EvictFirst => {
375 push_directive(tokens, "L2::evict_first");
376 }
377 Level2EvictionPriority::L2EvictLast => {
378 push_directive(tokens, "L2::evict_last");
379 }
380 }
381 }
382 if let Some(level_cache_hint_15) = self.level_cache_hint.as_ref() {
383 match level_cache_hint_15 {
384 LevelCacheHint::L2CacheHint => {
385 push_directive(tokens, "L2::cache_hint");
386 }
387 }
388 }
389 if let Some(level_prefetch_size_16) = self.level_prefetch_size.as_ref() {
390 match level_prefetch_size_16 {
391 LevelPrefetchSize::L2128b => {
392 push_directive(tokens, "L2::128B");
393 }
394 LevelPrefetchSize::L2256b => {
395 push_directive(tokens, "L2::256B");
396 }
397 LevelPrefetchSize::L264b => {
398 push_directive(tokens, "L2::64B");
399 }
400 }
401 }
402 match &self.vec {
403 Vec::V2 => {
404 push_directive(tokens, "v2");
405 }
406 Vec::V4 => {
407 push_directive(tokens, "v4");
408 }
409 Vec::V8 => {
410 push_directive(tokens, "v8");
411 }
412 }
413 match &self.type_ {
414 Type::B128 => {
415 push_directive(tokens, "b128");
416 }
417 Type::B16 => {
418 push_directive(tokens, "b16");
419 }
420 Type::B32 => {
421 push_directive(tokens, "b32");
422 }
423 Type::B64 => {
424 push_directive(tokens, "b64");
425 }
426 Type::U16 => {
427 push_directive(tokens, "u16");
428 }
429 Type::U32 => {
430 push_directive(tokens, "u32");
431 }
432 Type::U64 => {
433 push_directive(tokens, "u64");
434 }
435 Type::S16 => {
436 push_directive(tokens, "s16");
437 }
438 Type::S32 => {
439 push_directive(tokens, "s32");
440 }
441 Type::S64 => {
442 push_directive(tokens, "s64");
443 }
444 Type::F32 => {
445 push_directive(tokens, "f32");
446 }
447 Type::F64 => {
448 push_directive(tokens, "f64");
449 }
450 Type::B8 => {
451 push_directive(tokens, "b8");
452 }
453 Type::U8 => {
454 push_directive(tokens, "u8");
455 }
456 Type::S8 => {
457 push_directive(tokens, "s8");
458 }
459 }
460 self.d.unparse_tokens(tokens);
461 tokens.push(PtxToken::Comma);
462 self.a.unparse_tokens(tokens);
463 if self.cache_policy.is_some() { tokens.push(PtxToken::Comma); }
464 if let Some(opt_17) = self.cache_policy.as_ref() {
465 opt_17.unparse_tokens(tokens);
466 }
467 tokens.push(PtxToken::Semicolon);
468 }
469 }
470
471}
472