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() {
116                tokens.push(PtxToken::Comma);
117            }
118            if let Some(opt_3) = self.cache_policy.as_ref() {
119                opt_3.unparse_tokens(tokens);
120            }
121            tokens.push(PtxToken::Semicolon);
122        }
123    }
124
125    impl PtxUnparser for LdGlobalCopNcLevelCacheHintLevelPrefetchSizeVecType {
126        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
127            push_opcode(tokens, "ld");
128            push_directive(tokens, "global");
129            if let Some(cop_4) = self.cop.as_ref() {
130                match cop_4 {
131                    Cop::Ca => {
132                        push_directive(tokens, "ca");
133                    }
134                    Cop::Cg => {
135                        push_directive(tokens, "cg");
136                    }
137                    Cop::Cs => {
138                        push_directive(tokens, "cs");
139                    }
140                }
141            }
142            push_directive(tokens, "nc");
143            if let Some(level_cache_hint_5) = self.level_cache_hint.as_ref() {
144                match level_cache_hint_5 {
145                    LevelCacheHint::L2CacheHint => {
146                        push_directive(tokens, "L2::cache_hint");
147                    }
148                }
149            }
150            if let Some(level_prefetch_size_6) = self.level_prefetch_size.as_ref() {
151                match level_prefetch_size_6 {
152                    LevelPrefetchSize::L2128b => {
153                        push_directive(tokens, "L2::128B");
154                    }
155                    LevelPrefetchSize::L2256b => {
156                        push_directive(tokens, "L2::256B");
157                    }
158                    LevelPrefetchSize::L264b => {
159                        push_directive(tokens, "L2::64B");
160                    }
161                }
162            }
163            match &self.vec {
164                Vec::V2 => {
165                    push_directive(tokens, "v2");
166                }
167                Vec::V4 => {
168                    push_directive(tokens, "v4");
169                }
170                Vec::V8 => {
171                    push_directive(tokens, "v8");
172                }
173            }
174            match &self.type_ {
175                Type::B128 => {
176                    push_directive(tokens, "b128");
177                }
178                Type::B16 => {
179                    push_directive(tokens, "b16");
180                }
181                Type::B32 => {
182                    push_directive(tokens, "b32");
183                }
184                Type::B64 => {
185                    push_directive(tokens, "b64");
186                }
187                Type::U16 => {
188                    push_directive(tokens, "u16");
189                }
190                Type::U32 => {
191                    push_directive(tokens, "u32");
192                }
193                Type::U64 => {
194                    push_directive(tokens, "u64");
195                }
196                Type::S16 => {
197                    push_directive(tokens, "s16");
198                }
199                Type::S32 => {
200                    push_directive(tokens, "s32");
201                }
202                Type::S64 => {
203                    push_directive(tokens, "s64");
204                }
205                Type::F32 => {
206                    push_directive(tokens, "f32");
207                }
208                Type::F64 => {
209                    push_directive(tokens, "f64");
210                }
211                Type::B8 => {
212                    push_directive(tokens, "b8");
213                }
214                Type::U8 => {
215                    push_directive(tokens, "u8");
216                }
217                Type::S8 => {
218                    push_directive(tokens, "s8");
219                }
220            }
221            self.d.unparse_tokens(tokens);
222            tokens.push(PtxToken::Comma);
223            self.a.unparse_tokens(tokens);
224            if self.cache_policy.is_some() {
225                tokens.push(PtxToken::Comma);
226            }
227            if let Some(opt_7) = self.cache_policy.as_ref() {
228                opt_7.unparse_tokens(tokens);
229            }
230            tokens.push(PtxToken::Semicolon);
231        }
232    }
233
234    impl PtxUnparser for LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeType {
235        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
236            push_opcode(tokens, "ld");
237                    push_directive(tokens, "global");
238                    push_directive(tokens, "nc");
239                    if let Some(level1_eviction_priority_8) = self.level1_eviction_priority.as_ref() {
240                            match level1_eviction_priority_8 {
241                                    Level1EvictionPriority::L1EvictUnchanged => {
242                                            push_directive(tokens, "L1::evict_unchanged");
243                                    }
244                                    Level1EvictionPriority::L1EvictNormal => {
245                                            push_directive(tokens, "L1::evict_normal");
246                                    }
247                                    Level1EvictionPriority::L1EvictFirst => {
248                                            push_directive(tokens, "L1::evict_first");
249                                    }
250                                    Level1EvictionPriority::L1NoAllocate => {
251                                            push_directive(tokens, "L1::no_allocate");
252                                    }
253                                    Level1EvictionPriority::L1EvictLast => {
254                                            push_directive(tokens, "L1::evict_last");
255                                    }
256                            }
257                    }
258                    if let Some(level2_eviction_priority_9) = self.level2_eviction_priority.as_ref() {
259                            match level2_eviction_priority_9 {
260                                    Level2EvictionPriority::L2EvictNormal => {
261                                            push_directive(tokens, "L2::evict_normal");
262                                    }
263                                    Level2EvictionPriority::L2EvictFirst => {
264                                            push_directive(tokens, "L2::evict_first");
265                                    }
266                                    Level2EvictionPriority::L2EvictLast => {
267                                            push_directive(tokens, "L2::evict_last");
268                                    }
269                            }
270                    }
271                    if let Some(level_cache_hint_10) = self.level_cache_hint.as_ref() {
272                            match level_cache_hint_10 {
273                                    LevelCacheHint::L2CacheHint => {
274                                            push_directive(tokens, "L2::cache_hint");
275                                    }
276                            }
277                    }
278                    if let Some(level_prefetch_size_11) = self.level_prefetch_size.as_ref() {
279                            match level_prefetch_size_11 {
280                                    LevelPrefetchSize::L2128b => {
281                                            push_directive(tokens, "L2::128B");
282                                    }
283                                    LevelPrefetchSize::L2256b => {
284                                            push_directive(tokens, "L2::256B");
285                                    }
286                                    LevelPrefetchSize::L264b => {
287                                            push_directive(tokens, "L2::64B");
288                                    }
289                            }
290                    }
291                    match &self.type_ {
292                            Type::B128 => {
293                                    push_directive(tokens, "b128");
294                            }
295                            Type::B16 => {
296                                    push_directive(tokens, "b16");
297                            }
298                            Type::B32 => {
299                                    push_directive(tokens, "b32");
300                            }
301                            Type::B64 => {
302                                    push_directive(tokens, "b64");
303                            }
304                            Type::U16 => {
305                                    push_directive(tokens, "u16");
306                            }
307                            Type::U32 => {
308                                    push_directive(tokens, "u32");
309                            }
310                            Type::U64 => {
311                                    push_directive(tokens, "u64");
312                            }
313                            Type::S16 => {
314                                    push_directive(tokens, "s16");
315                            }
316                            Type::S32 => {
317                                    push_directive(tokens, "s32");
318                            }
319                            Type::S64 => {
320                                    push_directive(tokens, "s64");
321                            }
322                            Type::F32 => {
323                                    push_directive(tokens, "f32");
324                            }
325                            Type::F64 => {
326                                    push_directive(tokens, "f64");
327                            }
328                            Type::B8 => {
329                                    push_directive(tokens, "b8");
330                            }
331                            Type::U8 => {
332                                    push_directive(tokens, "u8");
333                            }
334                            Type::S8 => {
335                                    push_directive(tokens, "s8");
336                            }
337                    }
338                    self.d.unparse_tokens(tokens);
339            tokens.push(PtxToken::Comma);
340                    self.a.unparse_tokens(tokens);
341            if self.cache_policy.is_some() { tokens.push(PtxToken::Comma); }
342                    if let Some(opt_12) = self.cache_policy.as_ref() {
343                        opt_12.unparse_tokens(tokens);
344                    }
345            tokens.push(PtxToken::Semicolon);
346        }
347    }
348
349    impl PtxUnparser for LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType {
350        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
351            push_opcode(tokens, "ld");
352                    push_directive(tokens, "global");
353                    push_directive(tokens, "nc");
354                    if let Some(level1_eviction_priority_13) = self.level1_eviction_priority.as_ref() {
355                            match level1_eviction_priority_13 {
356                                    Level1EvictionPriority::L1EvictUnchanged => {
357                                            push_directive(tokens, "L1::evict_unchanged");
358                                    }
359                                    Level1EvictionPriority::L1EvictNormal => {
360                                            push_directive(tokens, "L1::evict_normal");
361                                    }
362                                    Level1EvictionPriority::L1EvictFirst => {
363                                            push_directive(tokens, "L1::evict_first");
364                                    }
365                                    Level1EvictionPriority::L1NoAllocate => {
366                                            push_directive(tokens, "L1::no_allocate");
367                                    }
368                                    Level1EvictionPriority::L1EvictLast => {
369                                            push_directive(tokens, "L1::evict_last");
370                                    }
371                            }
372                    }
373                    if let Some(level2_eviction_priority_14) = self.level2_eviction_priority.as_ref() {
374                            match level2_eviction_priority_14 {
375                                    Level2EvictionPriority::L2EvictNormal => {
376                                            push_directive(tokens, "L2::evict_normal");
377                                    }
378                                    Level2EvictionPriority::L2EvictFirst => {
379                                            push_directive(tokens, "L2::evict_first");
380                                    }
381                                    Level2EvictionPriority::L2EvictLast => {
382                                            push_directive(tokens, "L2::evict_last");
383                                    }
384                            }
385                    }
386                    if let Some(level_cache_hint_15) = self.level_cache_hint.as_ref() {
387                            match level_cache_hint_15 {
388                                    LevelCacheHint::L2CacheHint => {
389                                            push_directive(tokens, "L2::cache_hint");
390                                    }
391                            }
392                    }
393                    if let Some(level_prefetch_size_16) = self.level_prefetch_size.as_ref() {
394                            match level_prefetch_size_16 {
395                                    LevelPrefetchSize::L2128b => {
396                                            push_directive(tokens, "L2::128B");
397                                    }
398                                    LevelPrefetchSize::L2256b => {
399                                            push_directive(tokens, "L2::256B");
400                                    }
401                                    LevelPrefetchSize::L264b => {
402                                            push_directive(tokens, "L2::64B");
403                                    }
404                            }
405                    }
406                    match &self.vec {
407                            Vec::V2 => {
408                                    push_directive(tokens, "v2");
409                            }
410                            Vec::V4 => {
411                                    push_directive(tokens, "v4");
412                            }
413                            Vec::V8 => {
414                                    push_directive(tokens, "v8");
415                            }
416                    }
417                    match &self.type_ {
418                            Type::B128 => {
419                                    push_directive(tokens, "b128");
420                            }
421                            Type::B16 => {
422                                    push_directive(tokens, "b16");
423                            }
424                            Type::B32 => {
425                                    push_directive(tokens, "b32");
426                            }
427                            Type::B64 => {
428                                    push_directive(tokens, "b64");
429                            }
430                            Type::U16 => {
431                                    push_directive(tokens, "u16");
432                            }
433                            Type::U32 => {
434                                    push_directive(tokens, "u32");
435                            }
436                            Type::U64 => {
437                                    push_directive(tokens, "u64");
438                            }
439                            Type::S16 => {
440                                    push_directive(tokens, "s16");
441                            }
442                            Type::S32 => {
443                                    push_directive(tokens, "s32");
444                            }
445                            Type::S64 => {
446                                    push_directive(tokens, "s64");
447                            }
448                            Type::F32 => {
449                                    push_directive(tokens, "f32");
450                            }
451                            Type::F64 => {
452                                    push_directive(tokens, "f64");
453                            }
454                            Type::B8 => {
455                                    push_directive(tokens, "b8");
456                            }
457                            Type::U8 => {
458                                    push_directive(tokens, "u8");
459                            }
460                            Type::S8 => {
461                                    push_directive(tokens, "s8");
462                            }
463                    }
464                    self.d.unparse_tokens(tokens);
465            tokens.push(PtxToken::Comma);
466                    self.a.unparse_tokens(tokens);
467            if self.cache_policy.is_some() { tokens.push(PtxToken::Comma); }
468                    if let Some(opt_17) = self.cache_policy.as_ref() {
469                        opt_17.unparse_tokens(tokens);
470                    }
471            tokens.push(PtxToken::Semicolon);
472        }
473    }
474}