ptx_parser/parser/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::parser::{
10    PtxParseError, PtxParser, PtxTokenStream, Span,
11    util::{
12        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
13        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
14    },
15};
16use crate::r#type::common::*;
17use crate::{alt, ok, seq_n};
18
19pub mod section_0 {
20    use super::*;
21    use crate::r#type::instruction::cp_async_bulk_prefetch::section_0::*;
22
23    // ============================================================================
24    // Generated enum parsers
25    // ============================================================================
26
27    impl PtxParser for LevelCacheHint {
28        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
29            alt!(map(string_p(".L2::cache_hint"), |_, _span| {
30                LevelCacheHint::L2CacheHint
31            }))
32        }
33    }
34
35    impl PtxParser for Src {
36        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
37            alt!(map(string_p(".global"), |_, _span| Src::Global))
38        }
39    }
40
41    impl PtxParser for CpAsyncBulkPrefetchL2SrcLevelCacheHint {
42        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
43            try_map(
44                seq_n!(
45                    string_p("cp"),
46                    string_p(".async"),
47                    string_p(".bulk"),
48                    string_p(".prefetch"),
49                    string_p(".L2"),
50                    Src::parse(),
51                    optional(LevelCacheHint::parse()),
52                    AddressOperand::parse(),
53                    comma_p(),
54                    GeneralOperand::parse(),
55                    map(
56                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
57                        |value, _| value.map(|(_, operand)| operand)
58                    ),
59                    semicolon_p()
60                ),
61                |(
62                    _,
63                    async_,
64                    bulk,
65                    prefetch,
66                    l2,
67                    src,
68                    level_cache_hint,
69                    srcmem,
70                    _,
71                    size,
72                    cache_policy,
73                    _,
74                ),
75                 span| {
76                    ok!(CpAsyncBulkPrefetchL2SrcLevelCacheHint {
77                        async_ = async_,
78                        bulk = bulk,
79                        prefetch = prefetch,
80                        l2 = l2,
81                        src = src,
82                        level_cache_hint = level_cache_hint,
83                        srcmem = srcmem,
84                        size = size,
85                        cache_policy = cache_policy,
86
87                    })
88                },
89            )
90        }
91    }
92}
ptx_parser/parser/instruction/cp_async_bulk_prefetch.rs

ptx_parser/parser/instruction/
cp_async_bulk_prefetch.rs
