ptx_parser/parser/instruction/
cp_async_bulk_prefetch_tensor.rs

1//! Original PTX specification:
2//!
3//! // global -> shared::cluster:
4//! cp.async.bulk.prefetch.tensor.dim.L2.src{.load_mode}{.level::cache_hint} [tensorMap, tensorCoords] {, im2colInfo } {, cache-policy};
5//! .src =                { .global };
6//! .dim =                { .1d, .2d, .3d, .4d, .5d };
7//! .load_mode =          { .tile, .tile::gather4, .im2col, .im2col::w, .im2col::w::128 };
8//! .level::cache_hint =  { .L2::cache_hint };
9
10#![allow(unused)]
11
12use crate::lexer::PtxToken;
13use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
14use crate::r#type::common::*;
15
16pub mod section_0 {
17    use super::*;
18    use crate::r#type::instruction::cp_async_bulk_prefetch_tensor::section_0::*;
19
20    // ============================================================================
21    // Generated enum parsers
22    // ============================================================================
23
24    impl PtxParser for Dim {
25        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
26            // Try _1d
27            {
28                let saved_pos = stream.position();
29                if stream.expect_string(".1d").is_ok() {
30                    return Ok(Dim::_1d);
31                }
32                stream.set_position(saved_pos);
33            }
34            let saved_pos = stream.position();
35            // Try _2d
36            {
37                let saved_pos = stream.position();
38                if stream.expect_string(".2d").is_ok() {
39                    return Ok(Dim::_2d);
40                }
41                stream.set_position(saved_pos);
42            }
43            stream.set_position(saved_pos);
44            let saved_pos = stream.position();
45            // Try _3d
46            {
47                let saved_pos = stream.position();
48                if stream.expect_string(".3d").is_ok() {
49                    return Ok(Dim::_3d);
50                }
51                stream.set_position(saved_pos);
52            }
53            stream.set_position(saved_pos);
54            let saved_pos = stream.position();
55            // Try _4d
56            {
57                let saved_pos = stream.position();
58                if stream.expect_string(".4d").is_ok() {
59                    return Ok(Dim::_4d);
60                }
61                stream.set_position(saved_pos);
62            }
63            stream.set_position(saved_pos);
64            let saved_pos = stream.position();
65            // Try _5d
66            {
67                let saved_pos = stream.position();
68                if stream.expect_string(".5d").is_ok() {
69                    return Ok(Dim::_5d);
70                }
71                stream.set_position(saved_pos);
72            }
73            stream.set_position(saved_pos);
74            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
75            let expected = &[".1d", ".2d", ".3d", ".4d", ".5d"];
76            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
77            Err(crate::parser::unexpected_value(span, expected, found))
78        }
79    }
80
81    impl PtxParser for LevelCacheHint {
82        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
83            // Try L2CacheHint
84            {
85                let saved_pos = stream.position();
86                if stream.expect_string(".L2::cache_hint").is_ok() {
87                    return Ok(LevelCacheHint::L2CacheHint);
88                }
89                stream.set_position(saved_pos);
90            }
91            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
92            let expected = &[".L2::cache_hint"];
93            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
94            Err(crate::parser::unexpected_value(span, expected, found))
95        }
96    }
97
98    impl PtxParser for LoadMode {
99        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
100            // Try Im2colW128
101            {
102                let saved_pos = stream.position();
103                if stream.expect_string(".im2col::w::128").is_ok() {
104                    return Ok(LoadMode::Im2colW128);
105                }
106                stream.set_position(saved_pos);
107            }
108            let saved_pos = stream.position();
109            // Try TileGather4
110            {
111                let saved_pos = stream.position();
112                if stream.expect_string(".tile::gather4").is_ok() {
113                    return Ok(LoadMode::TileGather4);
114                }
115                stream.set_position(saved_pos);
116            }
117            stream.set_position(saved_pos);
118            let saved_pos = stream.position();
119            // Try Im2colW
120            {
121                let saved_pos = stream.position();
122                if stream.expect_string(".im2col::w").is_ok() {
123                    return Ok(LoadMode::Im2colW);
124                }
125                stream.set_position(saved_pos);
126            }
127            stream.set_position(saved_pos);
128            let saved_pos = stream.position();
129            // Try Im2col
130            {
131                let saved_pos = stream.position();
132                if stream.expect_string(".im2col").is_ok() {
133                    return Ok(LoadMode::Im2col);
134                }
135                stream.set_position(saved_pos);
136            }
137            stream.set_position(saved_pos);
138            let saved_pos = stream.position();
139            // Try Tile
140            {
141                let saved_pos = stream.position();
142                if stream.expect_string(".tile").is_ok() {
143                    return Ok(LoadMode::Tile);
144                }
145                stream.set_position(saved_pos);
146            }
147            stream.set_position(saved_pos);
148            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
149            let expected = &[".im2col::w::128", ".tile::gather4", ".im2col::w", ".im2col", ".tile"];
150            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
151            Err(crate::parser::unexpected_value(span, expected, found))
152        }
153    }
154
155    impl PtxParser for Src {
156        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
157            // Try Global
158            {
159                let saved_pos = stream.position();
160                if stream.expect_string(".global").is_ok() {
161                    return Ok(Src::Global);
162                }
163                stream.set_position(saved_pos);
164            }
165            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
166            let expected = &[".global"];
167            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
168            Err(crate::parser::unexpected_value(span, expected, found))
169        }
170    }
171
172    impl PtxParser for CpAsyncBulkPrefetchTensorDimL2SrcLoadModeLevelCacheHint {
173        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
174            stream.expect_string("cp")?;
175            stream.expect_string(".async")?;
176            let async_ = ();
177            stream.expect_complete()?;
178            stream.expect_string(".bulk")?;
179            let bulk = ();
180            stream.expect_complete()?;
181            stream.expect_string(".prefetch")?;
182            let prefetch = ();
183            stream.expect_complete()?;
184            stream.expect_string(".tensor")?;
185            let tensor = ();
186            stream.expect_complete()?;
187            let dim = Dim::parse(stream)?;
188            stream.expect_complete()?;
189            stream.expect_string(".L2")?;
190            let l2 = ();
191            stream.expect_complete()?;
192            let src = Src::parse(stream)?;
193            stream.expect_complete()?;
194            let saved_pos = stream.position();
195            let load_mode = match LoadMode::parse(stream) {
196                Ok(val) => Some(val),
197                Err(_) => {
198                    stream.set_position(saved_pos);
199                    None
200                }
201            };
202            stream.expect_complete()?;
203            let saved_pos = stream.position();
204            let level_cache_hint = match LevelCacheHint::parse(stream) {
205                Ok(val) => Some(val),
206                Err(_) => {
207                    stream.set_position(saved_pos);
208                    None
209                }
210            };
211            stream.expect_complete()?;
212            let tensormap = TexHandler2::parse(stream)?;
213            stream.expect_complete()?;
214            let saved_pos = stream.position();
215            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
216            if !has_comma {
217                stream.set_position(saved_pos);
218            }
219            let saved_pos = stream.position();
220            let im2colinfo = match GeneralOperand::parse(stream) {
221                Ok(val) => Some(val),
222                Err(_) => {
223                    stream.set_position(saved_pos);
224                    None
225                }
226            };
227            stream.expect_complete()?;
228            let saved_pos = stream.position();
229            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
230            if !has_comma {
231                stream.set_position(saved_pos);
232            }
233            let saved_pos = stream.position();
234            let cache_policy = match GeneralOperand::parse(stream) {
235                Ok(val) => Some(val),
236                Err(_) => {
237                    stream.set_position(saved_pos);
238                    None
239                }
240            };
241            stream.expect_complete()?;
242            stream.expect_complete()?;
243            stream.expect(&PtxToken::Semicolon)?;
244            Ok(CpAsyncBulkPrefetchTensorDimL2SrcLoadModeLevelCacheHint {
245                async_,
246                bulk,
247                prefetch,
248                tensor,
249                dim,
250                l2,
251                src,
252                load_mode,
253                level_cache_hint,
254                tensormap,
255                im2colinfo,
256                cache_policy,
257            })
258        }
259    }
260
261
262}
263
ptx_parser/parser/instruction/cp_async_bulk_prefetch_tensor.rs

ptx_parser/parser/instruction/
cp_async_bulk_prefetch_tensor.rs
