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::lexer::PtxToken;
15use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
16use crate::r#type::common::*;
17
18pub mod section_0 {
19    use super::*;
20    use crate::r#type::instruction::cp_async::section_0::*;
21
22    // ============================================================================
23    // Generated enum parsers
24    // ============================================================================
25
26    impl PtxParser for CpSize {
27        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
28            // Try _16
29            {
30                let saved_pos = stream.position();
31                if stream.expect_string("16").is_ok() {
32                    return Ok(CpSize::_16);
33                }
34                stream.set_position(saved_pos);
35            }
36            let saved_pos = stream.position();
37            // Try _4
38            {
39                let saved_pos = stream.position();
40                if stream.expect_string("4").is_ok() {
41                    return Ok(CpSize::_4);
42                }
43                stream.set_position(saved_pos);
44            }
45            stream.set_position(saved_pos);
46            let saved_pos = stream.position();
47            // Try _8
48            {
49                let saved_pos = stream.position();
50                if stream.expect_string("8").is_ok() {
51                    return Ok(CpSize::_8);
52                }
53                stream.set_position(saved_pos);
54            }
55            stream.set_position(saved_pos);
56            let span = stream
57                .peek()
58                .map(|(_, s)| s.clone())
59                .unwrap_or(Span { start: 0, end: 0 });
60            let expected = &["16", "4", "8"];
61            let found = stream
62                .peek()
63                .map(|(t, _)| format!("{:?}", t))
64                .unwrap_or_else(|_| "<end of input>".to_string());
65            Err(crate::parser::unexpected_value(span, expected, found))
66        }
67    }
68
69    impl PtxParser for LevelCacheHint {
70        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
71            // Try L2CacheHint
72            {
73                let saved_pos = stream.position();
74                if stream.expect_string(".L2::cache_hint").is_ok() {
75                    return Ok(LevelCacheHint::L2CacheHint);
76                }
77                stream.set_position(saved_pos);
78            }
79            let span = stream
80                .peek()
81                .map(|(_, s)| s.clone())
82                .unwrap_or(Span { start: 0, end: 0 });
83            let expected = &[".L2::cache_hint"];
84            let found = stream
85                .peek()
86                .map(|(t, _)| format!("{:?}", t))
87                .unwrap_or_else(|_| "<end of input>".to_string());
88            Err(crate::parser::unexpected_value(span, expected, found))
89        }
90    }
91
92    impl PtxParser for LevelPrefetchSize {
93        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
94            // Try L2128b
95            {
96                let saved_pos = stream.position();
97                if stream.expect_string(".L2::128B").is_ok() {
98                    return Ok(LevelPrefetchSize::L2128b);
99                }
100                stream.set_position(saved_pos);
101            }
102            let saved_pos = stream.position();
103            // Try L2256b
104            {
105                let saved_pos = stream.position();
106                if stream.expect_string(".L2::256B").is_ok() {
107                    return Ok(LevelPrefetchSize::L2256b);
108                }
109                stream.set_position(saved_pos);
110            }
111            stream.set_position(saved_pos);
112            let saved_pos = stream.position();
113            // Try L264b
114            {
115                let saved_pos = stream.position();
116                if stream.expect_string(".L2::64B").is_ok() {
117                    return Ok(LevelPrefetchSize::L264b);
118                }
119                stream.set_position(saved_pos);
120            }
121            stream.set_position(saved_pos);
122            let span = stream
123                .peek()
124                .map(|(_, s)| s.clone())
125                .unwrap_or(Span { start: 0, end: 0 });
126            let expected = &[".L2::128B", ".L2::256B", ".L2::64B"];
127            let found = stream
128                .peek()
129                .map(|(t, _)| format!("{:?}", t))
130                .unwrap_or_else(|_| "<end of input>".to_string());
131            Err(crate::parser::unexpected_value(span, expected, found))
132        }
133    }
134
135    impl PtxParser for State {
136        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
137            // Try SharedCta
138            {
139                let saved_pos = stream.position();
140                if stream.expect_string(".shared::cta").is_ok() {
141                    return Ok(State::SharedCta);
142                }
143                stream.set_position(saved_pos);
144            }
145            let saved_pos = stream.position();
146            // Try Shared
147            {
148                let saved_pos = stream.position();
149                if stream.expect_string(".shared").is_ok() {
150                    return Ok(State::Shared);
151                }
152                stream.set_position(saved_pos);
153            }
154            stream.set_position(saved_pos);
155            let span = stream
156                .peek()
157                .map(|(_, s)| s.clone())
158                .unwrap_or(Span { start: 0, end: 0 });
159            let expected = &[".shared::cta", ".shared"];
160            let found = stream
161                .peek()
162                .map(|(t, _)| format!("{:?}", t))
163                .unwrap_or_else(|_| "<end of input>".to_string());
164            Err(crate::parser::unexpected_value(span, expected, found))
165        }
166    }
167
168    impl PtxParser for CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize {
169        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
170            stream.expect_string("cp")?;
171            stream.expect_string(".async")?;
172            let async_ = ();
173            stream.expect_complete()?;
174            stream.expect_string(".ca")?;
175            let ca = ();
176            stream.expect_complete()?;
177            let state = State::parse(stream)?;
178            stream.expect_complete()?;
179            stream.expect_string(".global")?;
180            let global = ();
181            stream.expect_complete()?;
182            let saved_pos = stream.position();
183            let level_cache_hint = match LevelCacheHint::parse(stream) {
184                Ok(val) => Some(val),
185                Err(_) => {
186                    stream.set_position(saved_pos);
187                    None
188                }
189            };
190            stream.expect_complete()?;
191            let saved_pos = stream.position();
192            let level_prefetch_size = match LevelPrefetchSize::parse(stream) {
193                Ok(val) => Some(val),
194                Err(_) => {
195                    stream.set_position(saved_pos);
196                    None
197                }
198            };
199            stream.expect_complete()?;
200            let dst = AddressOperand::parse(stream)?;
201            stream.expect_complete()?;
202            stream.expect(&PtxToken::Comma)?;
203            let src = AddressOperand::parse(stream)?;
204            stream.expect_complete()?;
205            stream.expect(&PtxToken::Comma)?;
206            let cp_size = CpSize::parse(stream)?;
207            stream.expect_complete()?;
208            let saved_pos = stream.position();
209            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
210            if !has_comma {
211                stream.set_position(saved_pos);
212            }
213            let saved_pos = stream.position();
214            let src_size = match GeneralOperand::parse(stream) {
215                Ok(val) => Some(val),
216                Err(_) => {
217                    stream.set_position(saved_pos);
218                    None
219                }
220            };
221            stream.expect_complete()?;
222            let saved_pos = stream.position();
223            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
224            if !has_comma {
225                stream.set_position(saved_pos);
226            }
227            let saved_pos = stream.position();
228            let cache_policy = match GeneralOperand::parse(stream) {
229                Ok(val) => Some(val),
230                Err(_) => {
231                    stream.set_position(saved_pos);
232                    None
233                }
234            };
235            stream.expect_complete()?;
236            stream.expect_complete()?;
237            stream.expect(&PtxToken::Semicolon)?;
238            Ok(CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize {
239                async_,
240                ca,
241                state,
242                global,
243                level_cache_hint,
244                level_prefetch_size,
245                dst,
246                src,
247                cp_size,
248                src_size,
249                cache_policy,
250            })
251        }
252    }
253
254    impl PtxParser for CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize {
255        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
256            stream.expect_string("cp")?;
257            stream.expect_string(".async")?;
258            let async_ = ();
259            stream.expect_complete()?;
260            stream.expect_string(".cg")?;
261            let cg = ();
262            stream.expect_complete()?;
263            let state = State::parse(stream)?;
264            stream.expect_complete()?;
265            stream.expect_string(".global")?;
266            let global = ();
267            stream.expect_complete()?;
268            let saved_pos = stream.position();
269            let level_cache_hint = match LevelCacheHint::parse(stream) {
270                Ok(val) => Some(val),
271                Err(_) => {
272                    stream.set_position(saved_pos);
273                    None
274                }
275            };
276            stream.expect_complete()?;
277            let saved_pos = stream.position();
278            let level_prefetch_size = match LevelPrefetchSize::parse(stream) {
279                Ok(val) => Some(val),
280                Err(_) => {
281                    stream.set_position(saved_pos);
282                    None
283                }
284            };
285            stream.expect_complete()?;
286            let dst = AddressOperand::parse(stream)?;
287            stream.expect_complete()?;
288            stream.expect(&PtxToken::Comma)?;
289            let src = AddressOperand::parse(stream)?;
290            stream.expect_complete()?;
291            stream.expect(&PtxToken::Comma)?;
292            stream.expect_string("16")?;
293            let imm_16 = ();
294            stream.expect_complete()?;
295            let saved_pos = stream.position();
296            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
297            if !has_comma {
298                stream.set_position(saved_pos);
299            }
300            let saved_pos = stream.position();
301            let src_size = match GeneralOperand::parse(stream) {
302                Ok(val) => Some(val),
303                Err(_) => {
304                    stream.set_position(saved_pos);
305                    None
306                }
307            };
308            stream.expect_complete()?;
309            let saved_pos = stream.position();
310            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
311            if !has_comma {
312                stream.set_position(saved_pos);
313            }
314            let saved_pos = stream.position();
315            let cache_policy = match GeneralOperand::parse(stream) {
316                Ok(val) => Some(val),
317                Err(_) => {
318                    stream.set_position(saved_pos);
319                    None
320                }
321            };
322            stream.expect_complete()?;
323            stream.expect_complete()?;
324            stream.expect(&PtxToken::Semicolon)?;
325            Ok(CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize {
326                async_,
327                cg,
328                state,
329                global,
330                level_cache_hint,
331                level_prefetch_size,
332                dst,
333                src,
334                imm_16,
335                src_size,
336                cache_policy,
337            })
338        }
339    }
340
341    impl PtxParser for CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize1 {
342        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
343            stream.expect_string("cp")?;
344            stream.expect_string(".async")?;
345            let async_ = ();
346            stream.expect_complete()?;
347            stream.expect_string(".ca")?;
348            let ca = ();
349            stream.expect_complete()?;
350            let state = State::parse(stream)?;
351            stream.expect_complete()?;
352            stream.expect_string(".global")?;
353            let global = ();
354            stream.expect_complete()?;
355            let saved_pos = stream.position();
356            let level_cache_hint = match LevelCacheHint::parse(stream) {
357                Ok(val) => Some(val),
358                Err(_) => {
359                    stream.set_position(saved_pos);
360                    None
361                }
362            };
363            stream.expect_complete()?;
364            let saved_pos = stream.position();
365            let level_prefetch_size = match LevelPrefetchSize::parse(stream) {
366                Ok(val) => Some(val),
367                Err(_) => {
368                    stream.set_position(saved_pos);
369                    None
370                }
371            };
372            stream.expect_complete()?;
373            let dst = AddressOperand::parse(stream)?;
374            stream.expect_complete()?;
375            stream.expect(&PtxToken::Comma)?;
376            let src = AddressOperand::parse(stream)?;
377            stream.expect_complete()?;
378            stream.expect(&PtxToken::Comma)?;
379            let cp_size = CpSize::parse(stream)?;
380            stream.expect_complete()?;
381            let saved_pos = stream.position();
382            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
383            if !has_comma {
384                stream.set_position(saved_pos);
385            }
386            let saved_pos = stream.position();
387            let ignore_src = match GeneralOperand::parse(stream) {
388                Ok(val) => Some(val),
389                Err(_) => {
390                    stream.set_position(saved_pos);
391                    None
392                }
393            };
394            stream.expect_complete()?;
395            let saved_pos = stream.position();
396            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
397            if !has_comma {
398                stream.set_position(saved_pos);
399            }
400            let saved_pos = stream.position();
401            let cache_policy = match GeneralOperand::parse(stream) {
402                Ok(val) => Some(val),
403                Err(_) => {
404                    stream.set_position(saved_pos);
405                    None
406                }
407            };
408            stream.expect_complete()?;
409            stream.expect_complete()?;
410            stream.expect(&PtxToken::Semicolon)?;
411            Ok(CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize1 {
412                async_,
413                ca,
414                state,
415                global,
416                level_cache_hint,
417                level_prefetch_size,
418                dst,
419                src,
420                cp_size,
421                ignore_src,
422                cache_policy,
423            })
424        }
425    }
426
427    impl PtxParser for CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize1 {
428        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
429            stream.expect_string("cp")?;
430            stream.expect_string(".async")?;
431            let async_ = ();
432            stream.expect_complete()?;
433            stream.expect_string(".cg")?;
434            let cg = ();
435            stream.expect_complete()?;
436            let state = State::parse(stream)?;
437            stream.expect_complete()?;
438            stream.expect_string(".global")?;
439            let global = ();
440            stream.expect_complete()?;
441            let saved_pos = stream.position();
442            let level_cache_hint = match LevelCacheHint::parse(stream) {
443                Ok(val) => Some(val),
444                Err(_) => {
445                    stream.set_position(saved_pos);
446                    None
447                }
448            };
449            stream.expect_complete()?;
450            let saved_pos = stream.position();
451            let level_prefetch_size = match LevelPrefetchSize::parse(stream) {
452                Ok(val) => Some(val),
453                Err(_) => {
454                    stream.set_position(saved_pos);
455                    None
456                }
457            };
458            stream.expect_complete()?;
459            let dst = AddressOperand::parse(stream)?;
460            stream.expect_complete()?;
461            stream.expect(&PtxToken::Comma)?;
462            let src = AddressOperand::parse(stream)?;
463            stream.expect_complete()?;
464            stream.expect(&PtxToken::Comma)?;
465            stream.expect_string("16")?;
466            let imm_16 = ();
467            stream.expect_complete()?;
468            let saved_pos = stream.position();
469            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
470            if !has_comma {
471                stream.set_position(saved_pos);
472            }
473            let saved_pos = stream.position();
474            let ignore_src = match GeneralOperand::parse(stream) {
475                Ok(val) => Some(val),
476                Err(_) => {
477                    stream.set_position(saved_pos);
478                    None
479                }
480            };
481            stream.expect_complete()?;
482            let saved_pos = stream.position();
483            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
484            if !has_comma {
485                stream.set_position(saved_pos);
486            }
487            let saved_pos = stream.position();
488            let cache_policy = match GeneralOperand::parse(stream) {
489                Ok(val) => Some(val),
490                Err(_) => {
491                    stream.set_position(saved_pos);
492                    None
493                }
494            };
495            stream.expect_complete()?;
496            stream.expect_complete()?;
497            stream.expect(&PtxToken::Semicolon)?;
498            Ok(CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize1 {
499                async_,
500                cg,
501                state,
502                global,
503                level_cache_hint,
504                level_prefetch_size,
505                dst,
506                src,
507                imm_16,
508                ignore_src,
509                cache_policy,
510            })
511        }
512    }
513}
ptx_parser/parser/instruction/cp_async.rs

ptx_parser/parser/instruction/
cp_async.rs
