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

ptx_parser/parser/instruction/
cp_async_bulk_prefetch.rs
