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

ptx_parser/parser/instruction/
cp_async.rs
