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

ptx_parser/unparser/instruction/
ld_global_nc.rs
