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            self.unparse_tokens_mode(tokens, false);
19        }
20        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
21            push_opcode(tokens, "cp");
22            push_directive(tokens, "async");
23            push_directive(tokens, "bulk");
24            push_directive(tokens, "prefetch");
25            push_directive(tokens, "L2");
26            match &self.src {
27                Src::Global => {
28                    push_directive(tokens, "global");
29                }
30            }
31            if let Some(level_cache_hint_0) = self.level_cache_hint.as_ref() {
32                match level_cache_hint_0 {
33                    LevelCacheHint::L2CacheHint => {
34                        push_directive(tokens, "L2::cache_hint");
35                    }
36                }
37            }
38            if spaced {
39                tokens.push(PtxToken::Space);
40            }
41            self.srcmem.unparse_tokens_mode(tokens, spaced);
42            tokens.push(PtxToken::Comma);
43            if spaced {
44                tokens.push(PtxToken::Space);
45            }
46            self.size.unparse_tokens_mode(tokens, spaced);
47            if self.cache_policy.is_some() {
48                tokens.push(PtxToken::Comma);
49            }
50            if let Some(opt_1) = self.cache_policy.as_ref() {
51                if spaced {
52                    tokens.push(PtxToken::Space);
53                }
54                opt_1.unparse_tokens_mode(tokens, spaced);
55            }
56            tokens.push(PtxToken::Semicolon);
57            if spaced {
58                tokens.push(PtxToken::Newline);
59            }
60        }
61    }
62}
ptx_parser/unparser/instruction/cp_async_bulk_prefetch.rs

ptx_parser/unparser/instruction/
cp_async_bulk_prefetch.rs
