ptx_parser/parser/instruction/
cp_async_bulk_prefetch_tensor.rs

1//! Original PTX specification:
2//!
3//! // global -> shared::cluster:
4//! cp.async.bulk.prefetch.tensor.dim.L2.src{.load_mode}{.level::cache_hint} [tensorMap, tensorCoords] {, im2colInfo } {, cache-policy};
5//! .src =                { .global };
6//! .dim =                { .1d, .2d, .3d, .4d, .5d };
7//! .load_mode =          { .tile, .tile::gather4, .im2col, .im2col::w, .im2col::w::128 };
8//! .level::cache_hint =  { .L2::cache_hint };
9
10#![allow(unused)]
11
12use crate::parser::{
13    PtxParseError, PtxParser, PtxTokenStream, Span,
14    util::{
15        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
16        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
17    },
18};
19use crate::r#type::common::*;
20use crate::{alt, ok, seq_n};
21
22pub mod section_0 {
23    use super::*;
24    use crate::r#type::instruction::cp_async_bulk_prefetch_tensor::section_0::*;
25
26    // ============================================================================
27    // Generated enum parsers
28    // ============================================================================
29
30    impl PtxParser for Dim {
31        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
32            alt!(
33                map(string_p(".1d"), |_, _span| Dim::_1d),
34                map(string_p(".2d"), |_, _span| Dim::_2d),
35                map(string_p(".3d"), |_, _span| Dim::_3d),
36                map(string_p(".4d"), |_, _span| Dim::_4d),
37                map(string_p(".5d"), |_, _span| Dim::_5d)
38            )
39        }
40    }
41
42    impl PtxParser for LevelCacheHint {
43        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
44            alt!(map(string_p(".L2::cache_hint"), |_, _span| {
45                LevelCacheHint::L2CacheHint
46            }))
47        }
48    }
49
50    impl PtxParser for LoadMode {
51        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
52            alt!(
53                map(string_p(".im2col::w::128"), |_, _span| LoadMode::Im2colW128),
54                map(string_p(".tile::gather4"), |_, _span| LoadMode::TileGather4),
55                map(string_p(".im2col::w"), |_, _span| LoadMode::Im2colW),
56                map(string_p(".im2col"), |_, _span| LoadMode::Im2col),
57                map(string_p(".tile"), |_, _span| LoadMode::Tile)
58            )
59        }
60    }
61
62    impl PtxParser for Src {
63        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
64            alt!(map(string_p(".global"), |_, _span| Src::Global))
65        }
66    }
67
68    impl PtxParser for CpAsyncBulkPrefetchTensorDimL2SrcLoadModeLevelCacheHint {
69        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
70            try_map(
71                seq_n!(
72                    string_p("cp"),
73                    string_p(".async"),
74                    string_p(".bulk"),
75                    string_p(".prefetch"),
76                    string_p(".tensor"),
77                    Dim::parse(),
78                    string_p(".L2"),
79                    Src::parse(),
80                    optional(LoadMode::parse()),
81                    optional(LevelCacheHint::parse()),
82                    TexHandler2::parse(),
83                    map(
84                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
85                        |value, _| value.map(|(_, operand)| operand)
86                    ),
87                    map(
88                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
89                        |value, _| value.map(|(_, operand)| operand)
90                    ),
91                    semicolon_p()
92                ),
93                |(
94                    _,
95                    async_,
96                    bulk,
97                    prefetch,
98                    tensor,
99                    dim,
100                    l2,
101                    src,
102                    load_mode,
103                    level_cache_hint,
104                    tensormap,
105                    im2colinfo,
106                    cache_policy,
107                    _,
108                ),
109                 span| {
110                    ok!(CpAsyncBulkPrefetchTensorDimL2SrcLoadModeLevelCacheHint {
111                        async_ = async_,
112                        bulk = bulk,
113                        prefetch = prefetch,
114                        tensor = tensor,
115                        dim = dim,
116                        l2 = l2,
117                        src = src,
118                        load_mode = load_mode,
119                        level_cache_hint = level_cache_hint,
120                        tensormap = tensormap,
121                        im2colinfo = im2colinfo,
122                        cache_policy = cache_policy,
123
124                    })
125                },
126            )
127        }
128    }
129}
ptx_parser/parser/instruction/cp_async_bulk_prefetch_tensor.rs

ptx_parser/parser/instruction/
cp_async_bulk_prefetch_tensor.rs
