ptx_parser/parser/instruction/
cp_async.rs

1//! Original PTX specification:
2//!
3//! cp.async.ca.state.global{.level::cache_hint}{.level::prefetch_size} [dst], [src], cp-size{, src-size}{, cache-policy};
4//! cp.async.cg.state.global{.level::cache_hint}{.level::prefetch_size} [dst], [src], 16{, src-size}{, cache-policy};
5//! cp.async.ca.state.global{.level::cache_hint}{.level::prefetch_size} [dst], [src], cp-size{, ignore-src}{, cache-policy} ;
6//! cp.async.cg.state.global{.level::cache_hint}{.level::prefetch_size} [dst], [src], 16{, ignore-src}{, cache-policy} ;
7//! .level::cache_hint =     { .L2::cache_hint };
8//! .level::prefetch_size =  { .L2::64B, .L2::128B, .L2::256B };
9//! cp-size = { 4, 8, 16 };
10//! .state = { .shared, .shared::cta}
11
12#![allow(unused)]
13
14use crate::parser::{
15    PtxParseError, PtxParser, PtxTokenStream, Span,
16    util::{
17        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
18        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
19    },
20};
21use crate::r#type::common::*;
22use crate::{alt, ok, seq_n};
23
24pub mod section_0 {
25    use super::*;
26    use crate::r#type::instruction::cp_async::section_0::*;
27
28    // ============================================================================
29    // Generated enum parsers
30    // ============================================================================
31
32    impl PtxParser for CpSize {
33        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
34            alt!(
35                map(string_p("16"), |_, _span| CpSize::_16),
36                map(string_p("4"), |_, _span| CpSize::_4),
37                map(string_p("8"), |_, _span| CpSize::_8)
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 LevelPrefetchSize {
51        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
52            alt!(
53                map(string_p(".L2::128B"), |_, _span| LevelPrefetchSize::L2128b),
54                map(string_p(".L2::256B"), |_, _span| LevelPrefetchSize::L2256b),
55                map(string_p(".L2::64B"), |_, _span| LevelPrefetchSize::L264b)
56            )
57        }
58    }
59
60    impl PtxParser for State {
61        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
62            alt!(
63                map(string_p(".shared::cta"), |_, _span| State::SharedCta),
64                map(string_p(".shared"), |_, _span| State::Shared)
65            )
66        }
67    }
68
69    impl PtxParser for CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize {
70        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
71            try_map(
72                seq_n!(
73                    string_p("cp"),
74                    string_p(".async"),
75                    string_p(".ca"),
76                    State::parse(),
77                    string_p(".global"),
78                    optional(LevelCacheHint::parse()),
79                    optional(LevelPrefetchSize::parse()),
80                    AddressOperand::parse(),
81                    comma_p(),
82                    AddressOperand::parse(),
83                    comma_p(),
84                    CpSize::parse(),
85                    map(
86                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
87                        |value, _| value.map(|(_, operand)| operand)
88                    ),
89                    map(
90                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
91                        |value, _| value.map(|(_, operand)| operand)
92                    ),
93                    semicolon_p()
94                ),
95                |(
96                    _,
97                    async_,
98                    ca,
99                    state,
100                    global,
101                    level_cache_hint,
102                    level_prefetch_size,
103                    dst,
104                    _,
105                    src,
106                    _,
107                    cp_size,
108                    src_size,
109                    cache_policy,
110                    _,
111                ),
112                 span| {
113                    ok!(CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize {
114                        async_ = async_,
115                        ca = ca,
116                        state = state,
117                        global = global,
118                        level_cache_hint = level_cache_hint,
119                        level_prefetch_size = level_prefetch_size,
120                        dst = dst,
121                        src = src,
122                        cp_size = cp_size,
123                        src_size = src_size,
124                        cache_policy = cache_policy,
125
126                    })
127                },
128            )
129        }
130    }
131
132    impl PtxParser for CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize {
133        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
134            try_map(
135                seq_n!(
136                    string_p("cp"),
137                    string_p(".async"),
138                    string_p(".cg"),
139                    State::parse(),
140                    string_p(".global"),
141                    optional(LevelCacheHint::parse()),
142                    optional(LevelPrefetchSize::parse()),
143                    AddressOperand::parse(),
144                    comma_p(),
145                    AddressOperand::parse(),
146                    comma_p(),
147                    map(string_p("16"), |_, _| ()),
148                    map(
149                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
150                        |value, _| value.map(|(_, operand)| operand)
151                    ),
152                    map(
153                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
154                        |value, _| value.map(|(_, operand)| operand)
155                    ),
156                    semicolon_p()
157                ),
158                |(
159                    _,
160                    async_,
161                    cg,
162                    state,
163                    global,
164                    level_cache_hint,
165                    level_prefetch_size,
166                    dst,
167                    _,
168                    src,
169                    _,
170                    imm_16,
171                    src_size,
172                    cache_policy,
173                    _,
174                ),
175                 span| {
176                    ok!(CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize {
177                        async_ = async_,
178                        cg = cg,
179                        state = state,
180                        global = global,
181                        level_cache_hint = level_cache_hint,
182                        level_prefetch_size = level_prefetch_size,
183                        dst = dst,
184                        src = src,
185                        imm_16 = imm_16,
186                        src_size = src_size,
187                        cache_policy = cache_policy,
188
189                    })
190                },
191            )
192        }
193    }
194
195    impl PtxParser for CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize1 {
196        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
197            try_map(
198                seq_n!(
199                    string_p("cp"),
200                    string_p(".async"),
201                    string_p(".ca"),
202                    State::parse(),
203                    string_p(".global"),
204                    optional(LevelCacheHint::parse()),
205                    optional(LevelPrefetchSize::parse()),
206                    AddressOperand::parse(),
207                    comma_p(),
208                    AddressOperand::parse(),
209                    comma_p(),
210                    CpSize::parse(),
211                    map(
212                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
213                        |value, _| value.map(|(_, operand)| operand)
214                    ),
215                    map(
216                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
217                        |value, _| value.map(|(_, operand)| operand)
218                    ),
219                    semicolon_p()
220                ),
221                |(
222                    _,
223                    async_,
224                    ca,
225                    state,
226                    global,
227                    level_cache_hint,
228                    level_prefetch_size,
229                    dst,
230                    _,
231                    src,
232                    _,
233                    cp_size,
234                    ignore_src,
235                    cache_policy,
236                    _,
237                ),
238                 span| {
239                    ok!(CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize1 {
240                        async_ = async_,
241                        ca = ca,
242                        state = state,
243                        global = global,
244                        level_cache_hint = level_cache_hint,
245                        level_prefetch_size = level_prefetch_size,
246                        dst = dst,
247                        src = src,
248                        cp_size = cp_size,
249                        ignore_src = ignore_src,
250                        cache_policy = cache_policy,
251
252                    })
253                },
254            )
255        }
256    }
257
258    impl PtxParser for CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize1 {
259        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
260            try_map(
261                seq_n!(
262                    string_p("cp"),
263                    string_p(".async"),
264                    string_p(".cg"),
265                    State::parse(),
266                    string_p(".global"),
267                    optional(LevelCacheHint::parse()),
268                    optional(LevelPrefetchSize::parse()),
269                    AddressOperand::parse(),
270                    comma_p(),
271                    AddressOperand::parse(),
272                    comma_p(),
273                    map(string_p("16"), |_, _| ()),
274                    map(
275                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
276                        |value, _| value.map(|(_, operand)| operand)
277                    ),
278                    map(
279                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
280                        |value, _| value.map(|(_, operand)| operand)
281                    ),
282                    semicolon_p()
283                ),
284                |(
285                    _,
286                    async_,
287                    cg,
288                    state,
289                    global,
290                    level_cache_hint,
291                    level_prefetch_size,
292                    dst,
293                    _,
294                    src,
295                    _,
296                    imm_16,
297                    ignore_src,
298                    cache_policy,
299                    _,
300                ),
301                 span| {
302                    ok!(CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize1 {
303                        async_ = async_,
304                        cg = cg,
305                        state = state,
306                        global = global,
307                        level_cache_hint = level_cache_hint,
308                        level_prefetch_size = level_prefetch_size,
309                        dst = dst,
310                        src = src,
311                        imm_16 = imm_16,
312                        ignore_src = ignore_src,
313                        cache_policy = cache_policy,
314
315                    })
316                },
317            )
318        }
319    }
320}
ptx_parser/parser/instruction/cp_async.rs

ptx_parser/parser/instruction/
cp_async.rs
