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::lexer::PtxToken;
10use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
11use crate::r#type::common::*;
12
13pub mod section_0 {
14    use super::*;
15    use crate::r#type::instruction::cp_async_bulk_prefetch::section_0::*;
16
17    // ============================================================================
18    // Generated enum parsers
19    // ============================================================================
20
21    impl PtxParser for LevelCacheHint {
22        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
23            // Try L2CacheHint
24            {
25                let saved_pos = stream.position();
26                if stream.expect_string(".L2::cache_hint").is_ok() {
27                    return Ok(LevelCacheHint::L2CacheHint);
28                }
29                stream.set_position(saved_pos);
30            }
31            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
32            let expected = &[".L2::cache_hint"];
33            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
34            Err(crate::parser::unexpected_value(span, expected, found))
35        }
36    }
37
38    impl PtxParser for Src {
39        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
40            // Try Global
41            {
42                let saved_pos = stream.position();
43                if stream.expect_string(".global").is_ok() {
44                    return Ok(Src::Global);
45                }
46                stream.set_position(saved_pos);
47            }
48            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
49            let expected = &[".global"];
50            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
51            Err(crate::parser::unexpected_value(span, expected, found))
52        }
53    }
54
55    impl PtxParser for CpAsyncBulkPrefetchL2SrcLevelCacheHint {
56        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
57            stream.expect_string("cp")?;
58            stream.expect_string(".async")?;
59            let async_ = ();
60            stream.expect_complete()?;
61            stream.expect_string(".bulk")?;
62            let bulk = ();
63            stream.expect_complete()?;
64            stream.expect_string(".prefetch")?;
65            let prefetch = ();
66            stream.expect_complete()?;
67            stream.expect_string(".L2")?;
68            let l2 = ();
69            stream.expect_complete()?;
70            let src = Src::parse(stream)?;
71            stream.expect_complete()?;
72            let saved_pos = stream.position();
73            let level_cache_hint = match LevelCacheHint::parse(stream) {
74                Ok(val) => Some(val),
75                Err(_) => {
76                    stream.set_position(saved_pos);
77                    None
78                }
79            };
80            stream.expect_complete()?;
81            let srcmem = AddressOperand::parse(stream)?;
82            stream.expect_complete()?;
83            stream.expect(&PtxToken::Comma)?;
84            let size = GeneralOperand::parse(stream)?;
85            stream.expect_complete()?;
86            let saved_pos = stream.position();
87            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
88            if !has_comma {
89                stream.set_position(saved_pos);
90            }
91            let saved_pos = stream.position();
92            let cache_policy = match GeneralOperand::parse(stream) {
93                Ok(val) => Some(val),
94                Err(_) => {
95                    stream.set_position(saved_pos);
96                    None
97                }
98            };
99            stream.expect_complete()?;
100            stream.expect_complete()?;
101            stream.expect(&PtxToken::Semicolon)?;
102            Ok(CpAsyncBulkPrefetchL2SrcLevelCacheHint {
103                async_,
104                bulk,
105                prefetch,
106                l2,
107                src,
108                level_cache_hint,
109                srcmem,
110                size,
111                cache_policy,
112            })
113        }
114    }
115
116
117}
118
ptx_parser/parser/instruction/cp_async_bulk_prefetch.rs

ptx_parser/parser/instruction/
cp_async_bulk_prefetch.rs
