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() {
39                tokens.push(PtxToken::Comma);
40            }
41            if let Some(opt_1) = self.cache_policy.as_ref() {
42                opt_1.unparse_tokens(tokens);
43            }
44            tokens.push(PtxToken::Semicolon);
45        }
46    }
47}
ptx_parser/unparser/instruction/cp_async_bulk_prefetch.rs

ptx_parser/unparser/instruction/
cp_async_bulk_prefetch.rs
