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
75                .peek()
76                .map(|(_, s)| s.clone())
77                .unwrap_or(Span { start: 0, end: 0 });
78            let expected = &[".1d", ".2d", ".3d", ".4d", ".5d"];
79            let found = stream
80                .peek()
81                .map(|(t, _)| format!("{:?}", t))
82                .unwrap_or_else(|_| "<end of input>".to_string());
83            Err(crate::parser::unexpected_value(span, expected, found))
84        }
85    }
86
87    impl PtxParser for LevelCacheHint {
88        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
89            // Try L2CacheHint
90            {
91                let saved_pos = stream.position();
92                if stream.expect_string(".L2::cache_hint").is_ok() {
93                    return Ok(LevelCacheHint::L2CacheHint);
94                }
95                stream.set_position(saved_pos);
96            }
97            let span = stream
98                .peek()
99                .map(|(_, s)| s.clone())
100                .unwrap_or(Span { start: 0, end: 0 });
101            let expected = &[".L2::cache_hint"];
102            let found = stream
103                .peek()
104                .map(|(t, _)| format!("{:?}", t))
105                .unwrap_or_else(|_| "<end of input>".to_string());
106            Err(crate::parser::unexpected_value(span, expected, found))
107        }
108    }
109
110    impl PtxParser for LoadMode {
111        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
112            // Try Im2colW128
113            {
114                let saved_pos = stream.position();
115                if stream.expect_string(".im2col::w::128").is_ok() {
116                    return Ok(LoadMode::Im2colW128);
117                }
118                stream.set_position(saved_pos);
119            }
120            let saved_pos = stream.position();
121            // Try TileGather4
122            {
123                let saved_pos = stream.position();
124                if stream.expect_string(".tile::gather4").is_ok() {
125                    return Ok(LoadMode::TileGather4);
126                }
127                stream.set_position(saved_pos);
128            }
129            stream.set_position(saved_pos);
130            let saved_pos = stream.position();
131            // Try Im2colW
132            {
133                let saved_pos = stream.position();
134                if stream.expect_string(".im2col::w").is_ok() {
135                    return Ok(LoadMode::Im2colW);
136                }
137                stream.set_position(saved_pos);
138            }
139            stream.set_position(saved_pos);
140            let saved_pos = stream.position();
141            // Try Im2col
142            {
143                let saved_pos = stream.position();
144                if stream.expect_string(".im2col").is_ok() {
145                    return Ok(LoadMode::Im2col);
146                }
147                stream.set_position(saved_pos);
148            }
149            stream.set_position(saved_pos);
150            let saved_pos = stream.position();
151            // Try Tile
152            {
153                let saved_pos = stream.position();
154                if stream.expect_string(".tile").is_ok() {
155                    return Ok(LoadMode::Tile);
156                }
157                stream.set_position(saved_pos);
158            }
159            stream.set_position(saved_pos);
160            let span = stream
161                .peek()
162                .map(|(_, s)| s.clone())
163                .unwrap_or(Span { start: 0, end: 0 });
164            let expected = &[
165                ".im2col::w::128",
166                ".tile::gather4",
167                ".im2col::w",
168                ".im2col",
169                ".tile",
170            ];
171            let found = stream
172                .peek()
173                .map(|(t, _)| format!("{:?}", t))
174                .unwrap_or_else(|_| "<end of input>".to_string());
175            Err(crate::parser::unexpected_value(span, expected, found))
176        }
177    }
178
179    impl PtxParser for Src {
180        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
181            // Try Global
182            {
183                let saved_pos = stream.position();
184                if stream.expect_string(".global").is_ok() {
185                    return Ok(Src::Global);
186                }
187                stream.set_position(saved_pos);
188            }
189            let span = stream
190                .peek()
191                .map(|(_, s)| s.clone())
192                .unwrap_or(Span { start: 0, end: 0 });
193            let expected = &[".global"];
194            let found = stream
195                .peek()
196                .map(|(t, _)| format!("{:?}", t))
197                .unwrap_or_else(|_| "<end of input>".to_string());
198            Err(crate::parser::unexpected_value(span, expected, found))
199        }
200    }
201
202    impl PtxParser for CpAsyncBulkPrefetchTensorDimL2SrcLoadModeLevelCacheHint {
203        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
204            stream.expect_string("cp")?;
205            stream.expect_string(".async")?;
206            let async_ = ();
207            stream.expect_complete()?;
208            stream.expect_string(".bulk")?;
209            let bulk = ();
210            stream.expect_complete()?;
211            stream.expect_string(".prefetch")?;
212            let prefetch = ();
213            stream.expect_complete()?;
214            stream.expect_string(".tensor")?;
215            let tensor = ();
216            stream.expect_complete()?;
217            let dim = Dim::parse(stream)?;
218            stream.expect_complete()?;
219            stream.expect_string(".L2")?;
220            let l2 = ();
221            stream.expect_complete()?;
222            let src = Src::parse(stream)?;
223            stream.expect_complete()?;
224            let saved_pos = stream.position();
225            let load_mode = match LoadMode::parse(stream) {
226                Ok(val) => Some(val),
227                Err(_) => {
228                    stream.set_position(saved_pos);
229                    None
230                }
231            };
232            stream.expect_complete()?;
233            let saved_pos = stream.position();
234            let level_cache_hint = match LevelCacheHint::parse(stream) {
235                Ok(val) => Some(val),
236                Err(_) => {
237                    stream.set_position(saved_pos);
238                    None
239                }
240            };
241            stream.expect_complete()?;
242            let tensormap = TexHandler2::parse(stream)?;
243            stream.expect_complete()?;
244            let saved_pos = stream.position();
245            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
246            if !has_comma {
247                stream.set_position(saved_pos);
248            }
249            let saved_pos = stream.position();
250            let im2colinfo = match GeneralOperand::parse(stream) {
251                Ok(val) => Some(val),
252                Err(_) => {
253                    stream.set_position(saved_pos);
254                    None
255                }
256            };
257            stream.expect_complete()?;
258            let saved_pos = stream.position();
259            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
260            if !has_comma {
261                stream.set_position(saved_pos);
262            }
263            let saved_pos = stream.position();
264            let cache_policy = match GeneralOperand::parse(stream) {
265                Ok(val) => Some(val),
266                Err(_) => {
267                    stream.set_position(saved_pos);
268                    None
269                }
270            };
271            stream.expect_complete()?;
272            stream.expect_complete()?;
273            stream.expect(&PtxToken::Semicolon)?;
274            Ok(CpAsyncBulkPrefetchTensorDimL2SrcLoadModeLevelCacheHint {
275                async_,
276                bulk,
277                prefetch,
278                tensor,
279                dim,
280                l2,
281                src,
282                load_mode,
283                level_cache_hint,
284                tensormap,
285                im2colinfo,
286                cache_policy,
287            })
288        }
289    }
290}
ptx_parser/parser/instruction/cp_async_bulk_prefetch_tensor.rs

ptx_parser/parser/instruction/
cp_async_bulk_prefetch_tensor.rs
