ptx_parser/unparser/instruction/
cp_async_bulk_prefetch.rs

1//! Original PTX specification:
2//!
3//! cp.async.bulk.prefetch.L2.src{.level::cache_hint}   [srcMem], size {, cache-policy};
4//! .src =                { .global };
5//! .level::cache_hint =  { .L2::cache_hint };
6
7#![allow(unused)]
8
9use crate::lexer::PtxToken;
10use crate::unparser::{PtxUnparser, common::*};
11
12pub mod section_0 {
13    use super::*;
14    use crate::r#type::instruction::cp_async_bulk_prefetch::section_0::*;
15
16    impl PtxUnparser for CpAsyncBulkPrefetchL2SrcLevelCacheHint {
17        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
18            push_opcode(tokens, "cp");
19                    push_directive(tokens, "async");
20                    push_directive(tokens, "bulk");
21                    push_directive(tokens, "prefetch");
22                    push_directive(tokens, "L2");
23                    match &self.src {
24                            Src::Global => {
25                                    push_directive(tokens, "global");
26                            }
27                    }
28                    if let Some(level_cache_hint_0) = self.level_cache_hint.as_ref() {
29                            match level_cache_hint_0 {
30                                    LevelCacheHint::L2CacheHint => {
31                                            push_directive(tokens, "L2::cache_hint");
32                                    }
33                            }
34                    }
35                    self.srcmem.unparse_tokens(tokens);
36            tokens.push(PtxToken::Comma);
37                    self.size.unparse_tokens(tokens);
38            if self.cache_policy.is_some() { tokens.push(PtxToken::Comma); }
39                    if let Some(opt_1) = self.cache_policy.as_ref() {
40                        opt_1.unparse_tokens(tokens);
41                    }
42            tokens.push(PtxToken::Semicolon);
43        }
44    }
45
46}
47