ptx_parser/parser/instruction/
ld_global_nc.rs

1//! Original PTX specification:
2//!
3//! ld.global{.cop}.nc{.level::cache_hint}{.level::prefetch_size}.type                 d, [a]{, cache-policy};
4//! ld.global{.cop}.nc{.level::cache_hint}{.level::prefetch_size}.vec.type             d, [a]{, cache-policy};
5//! ld.global.nc{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.level::prefetch_size}.type      d, [a]{, cache-policy};
6//! ld.global.nc{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.level::prefetch_size}.vec.type  d, [a]{, cache-policy};
7//! .cop  =                     { .ca, .cg, .cs };     // cache operation
8//! .level1::eviction_priority = { .L1::evict_normal, .L1::evict_unchanged, .L1::evict_first, .L1::evict_last, .L1::no_allocate};
9//! .level2::eviction_priority = {.L2::evict_normal, .L2::evict_first, .L2::evict_last};
10//! .level::cache_hint =        { .L2::cache_hint };
11//! .level::prefetch_size =     { .L2::64B, .L2::128B, .L2::256B };
12//! .vec  =                     { .v2, .v4, .v8 };
13//! .type =                     { .b8, .b16, .b32, .b64, .b128,
14//! .u8, .u16, .u32, .u64,
15//! .s8, .s16, .s32, .s64,
16//! .f32, .f64 };
17
18#![allow(unused)]
19
20use crate::parser::{
21    PtxParseError, PtxParser, PtxTokenStream, Span,
22    util::{
23        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
24        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
25    },
26};
27use crate::r#type::common::*;
28use crate::{alt, ok, seq_n};
29
30pub mod section_0 {
31    use super::*;
32    use crate::r#type::instruction::ld_global_nc::section_0::*;
33
34    // ============================================================================
35    // Generated enum parsers
36    // ============================================================================
37
38    impl PtxParser for Cop {
39        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
40            alt!(
41                map(string_p(".ca"), |_, _span| Cop::Ca),
42                map(string_p(".cg"), |_, _span| Cop::Cg),
43                map(string_p(".cs"), |_, _span| Cop::Cs)
44            )
45        }
46    }
47
48    impl PtxParser for Level1EvictionPriority {
49        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
50            alt!(
51                map(string_p(".L1::evict_unchanged"), |_, _span| {
52                    Level1EvictionPriority::L1EvictUnchanged
53                }),
54                map(string_p(".L1::evict_normal"), |_, _span| {
55                    Level1EvictionPriority::L1EvictNormal
56                }),
57                map(string_p(".L1::evict_first"), |_, _span| {
58                    Level1EvictionPriority::L1EvictFirst
59                }),
60                map(string_p(".L1::no_allocate"), |_, _span| {
61                    Level1EvictionPriority::L1NoAllocate
62                }),
63                map(string_p(".L1::evict_last"), |_, _span| {
64                    Level1EvictionPriority::L1EvictLast
65                })
66            )
67        }
68    }
69
70    impl PtxParser for Level2EvictionPriority {
71        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
72            alt!(
73                map(string_p(".L2::evict_normal"), |_, _span| {
74                    Level2EvictionPriority::L2EvictNormal
75                }),
76                map(string_p(".L2::evict_first"), |_, _span| {
77                    Level2EvictionPriority::L2EvictFirst
78                }),
79                map(string_p(".L2::evict_last"), |_, _span| {
80                    Level2EvictionPriority::L2EvictLast
81                })
82            )
83        }
84    }
85
86    impl PtxParser for LevelCacheHint {
87        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
88            alt!(map(string_p(".L2::cache_hint"), |_, _span| {
89                LevelCacheHint::L2CacheHint
90            }))
91        }
92    }
93
94    impl PtxParser for LevelPrefetchSize {
95        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
96            alt!(
97                map(string_p(".L2::128B"), |_, _span| LevelPrefetchSize::L2128b),
98                map(string_p(".L2::256B"), |_, _span| LevelPrefetchSize::L2256b),
99                map(string_p(".L2::64B"), |_, _span| LevelPrefetchSize::L264b)
100            )
101        }
102    }
103
104    impl PtxParser for Type {
105        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
106            alt!(
107                map(string_p(".b128"), |_, _span| Type::B128),
108                map(string_p(".b16"), |_, _span| Type::B16),
109                map(string_p(".b32"), |_, _span| Type::B32),
110                map(string_p(".b64"), |_, _span| Type::B64),
111                map(string_p(".u16"), |_, _span| Type::U16),
112                map(string_p(".u32"), |_, _span| Type::U32),
113                map(string_p(".u64"), |_, _span| Type::U64),
114                map(string_p(".s16"), |_, _span| Type::S16),
115                map(string_p(".s32"), |_, _span| Type::S32),
116                map(string_p(".s64"), |_, _span| Type::S64),
117                map(string_p(".f32"), |_, _span| Type::F32),
118                map(string_p(".f64"), |_, _span| Type::F64),
119                map(string_p(".b8"), |_, _span| Type::B8),
120                map(string_p(".u8"), |_, _span| Type::U8),
121                map(string_p(".s8"), |_, _span| Type::S8)
122            )
123        }
124    }
125
126    impl PtxParser for Vec {
127        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
128            alt!(
129                map(string_p(".v2"), |_, _span| Vec::V2),
130                map(string_p(".v4"), |_, _span| Vec::V4),
131                map(string_p(".v8"), |_, _span| Vec::V8)
132            )
133        }
134    }
135
136    impl PtxParser for LdGlobalCopNcLevelCacheHintLevelPrefetchSizeType {
137        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
138            try_map(
139                seq_n!(
140                    string_p("ld"),
141                    string_p(".global"),
142                    optional(Cop::parse()),
143                    string_p(".nc"),
144                    optional(LevelCacheHint::parse()),
145                    optional(LevelPrefetchSize::parse()),
146                    Type::parse(),
147                    GeneralOperand::parse(),
148                    comma_p(),
149                    AddressOperand::parse(),
150                    map(
151                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
152                        |value, _| value.map(|(_, operand)| operand)
153                    ),
154                    semicolon_p()
155                ),
156                |(
157                    _,
158                    global,
159                    cop,
160                    nc,
161                    level_cache_hint,
162                    level_prefetch_size,
163                    type_,
164                    d,
165                    _,
166                    a,
167                    cache_policy,
168                    _,
169                ),
170                 span| {
171                    ok!(LdGlobalCopNcLevelCacheHintLevelPrefetchSizeType {
172                        global = global,
173                        cop = cop,
174                        nc = nc,
175                        level_cache_hint = level_cache_hint,
176                        level_prefetch_size = level_prefetch_size,
177                        type_ = type_,
178                        d = d,
179                        a = a,
180                        cache_policy = cache_policy,
181
182                    })
183                },
184            )
185        }
186    }
187
188    impl PtxParser for LdGlobalCopNcLevelCacheHintLevelPrefetchSizeVecType {
189        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
190            try_map(
191                seq_n!(
192                    string_p("ld"),
193                    string_p(".global"),
194                    optional(Cop::parse()),
195                    string_p(".nc"),
196                    optional(LevelCacheHint::parse()),
197                    optional(LevelPrefetchSize::parse()),
198                    Vec::parse(),
199                    Type::parse(),
200                    GeneralOperand::parse(),
201                    comma_p(),
202                    AddressOperand::parse(),
203                    map(
204                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
205                        |value, _| value.map(|(_, operand)| operand)
206                    ),
207                    semicolon_p()
208                ),
209                |(
210                    _,
211                    global,
212                    cop,
213                    nc,
214                    level_cache_hint,
215                    level_prefetch_size,
216                    vec,
217                    type_,
218                    d,
219                    _,
220                    a,
221                    cache_policy,
222                    _,
223                ),
224                 span| {
225                    ok!(LdGlobalCopNcLevelCacheHintLevelPrefetchSizeVecType {
226                        global = global,
227                        cop = cop,
228                        nc = nc,
229                        level_cache_hint = level_cache_hint,
230                        level_prefetch_size = level_prefetch_size,
231                        vec = vec,
232                        type_ = type_,
233                        d = d,
234                        a = a,
235                        cache_policy = cache_policy,
236
237                    })
238                },
239            )
240        }
241    }
242
243    impl PtxParser for LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeType {
244        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
245            try_map(
246                seq_n!(
247                    string_p("ld"),
248                    string_p(".global"),
249                    string_p(".nc"),
250                    optional(Level1EvictionPriority::parse()),
251                    optional(Level2EvictionPriority::parse()),
252                    optional(LevelCacheHint::parse()),
253                    optional(LevelPrefetchSize::parse()),
254                    Type::parse(),
255                    GeneralOperand::parse(),
256                    comma_p(),
257                    AddressOperand::parse(),
258                    map(optional(seq_n!(comma_p(), GeneralOperand::parse())), |value, _| value.map(|(_, operand)| operand)),
259                    semicolon_p()
260                ),
261                |(_, global, nc, level1_eviction_priority, level2_eviction_priority, level_cache_hint, level_prefetch_size, type_, d, _, a, cache_policy, _), span| {
262                    ok!(LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeType {
263                        global = global,
264                        nc = nc,
265                        level1_eviction_priority = level1_eviction_priority,
266                        level2_eviction_priority = level2_eviction_priority,
267                        level_cache_hint = level_cache_hint,
268                        level_prefetch_size = level_prefetch_size,
269                        type_ = type_,
270                        d = d,
271                        a = a,
272                        cache_policy = cache_policy,
273
274                    })
275                },
276            )
277        }
278    }
279
280    impl PtxParser for LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType {
281        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
282            try_map(
283                seq_n!(
284                    string_p("ld"),
285                    string_p(".global"),
286                    string_p(".nc"),
287                    optional(Level1EvictionPriority::parse()),
288                    optional(Level2EvictionPriority::parse()),
289                    optional(LevelCacheHint::parse()),
290                    optional(LevelPrefetchSize::parse()),
291                    Vec::parse(),
292                    Type::parse(),
293                    GeneralOperand::parse(),
294                    comma_p(),
295                    AddressOperand::parse(),
296                    map(optional(seq_n!(comma_p(), GeneralOperand::parse())), |value, _| value.map(|(_, operand)| operand)),
297                    semicolon_p()
298                ),
299                |(_, global, nc, level1_eviction_priority, level2_eviction_priority, level_cache_hint, level_prefetch_size, vec, type_, d, _, a, cache_policy, _), span| {
300                    ok!(LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType {
301                        global = global,
302                        nc = nc,
303                        level1_eviction_priority = level1_eviction_priority,
304                        level2_eviction_priority = level2_eviction_priority,
305                        level_cache_hint = level_cache_hint,
306                        level_prefetch_size = level_prefetch_size,
307                        vec = vec,
308                        type_ = type_,
309                        d = d,
310                        a = a,
311                        cache_policy = cache_policy,
312
313                    })
314                },
315            )
316        }
317    }
318}
ptx_parser/parser/instruction/ld_global_nc.rs

ptx_parser/parser/instruction/
ld_global_nc.rs
