ptx_parser/parser/instruction/
ld_global_nc.rs

1//! Original PTX specification:
2//!
3//! ld.global{.cop}.nc{.level::cache_hint}{.level::prefetch_size}.type                 d, [a]{, cache-policy};
4//! ld.global{.cop}.nc{.level::cache_hint}{.level::prefetch_size}.vec.type             d, [a]{, cache-policy};
5//! ld.global.nc{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.level::prefetch_size}.type      d, [a]{, cache-policy};
6//! ld.global.nc{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.level::prefetch_size}.vec.type  d, [a]{, cache-policy};
7//! .cop  =                     { .ca, .cg, .cs };     // cache operation
8//! .level1::eviction_priority = { .L1::evict_normal, .L1::evict_unchanged, .L1::evict_first, .L1::evict_last, .L1::no_allocate};
9//! .level2::eviction_priority = {.L2::evict_normal, .L2::evict_first, .L2::evict_last};
10//! .level::cache_hint =        { .L2::cache_hint };
11//! .level::prefetch_size =     { .L2::64B, .L2::128B, .L2::256B };
12//! .vec  =                     { .v2, .v4, .v8 };
13//! .type =                     { .b8, .b16, .b32, .b64, .b128,
14//! .u8, .u16, .u32, .u64,
15//! .s8, .s16, .s32, .s64,
16//! .f32, .f64 };
17
18#![allow(unused)]
19
20use crate::lexer::PtxToken;
21use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
22use crate::r#type::common::*;
23
24pub mod section_0 {
25    use super::*;
26    use crate::r#type::instruction::ld_global_nc::section_0::*;
27
28    // ============================================================================
29    // Generated enum parsers
30    // ============================================================================
31
32    impl PtxParser for Cop {
33        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
34            // Try Ca
35            {
36                let saved_pos = stream.position();
37                if stream.expect_string(".ca").is_ok() {
38                    return Ok(Cop::Ca);
39                }
40                stream.set_position(saved_pos);
41            }
42            let saved_pos = stream.position();
43            // Try Cg
44            {
45                let saved_pos = stream.position();
46                if stream.expect_string(".cg").is_ok() {
47                    return Ok(Cop::Cg);
48                }
49                stream.set_position(saved_pos);
50            }
51            stream.set_position(saved_pos);
52            let saved_pos = stream.position();
53            // Try Cs
54            {
55                let saved_pos = stream.position();
56                if stream.expect_string(".cs").is_ok() {
57                    return Ok(Cop::Cs);
58                }
59                stream.set_position(saved_pos);
60            }
61            stream.set_position(saved_pos);
62            let span = stream
63                .peek()
64                .map(|(_, s)| s.clone())
65                .unwrap_or(Span { start: 0, end: 0 });
66            let expected = &[".ca", ".cg", ".cs"];
67            let found = stream
68                .peek()
69                .map(|(t, _)| format!("{:?}", t))
70                .unwrap_or_else(|_| "<end of input>".to_string());
71            Err(crate::parser::unexpected_value(span, expected, found))
72        }
73    }
74
75    impl PtxParser for Level1EvictionPriority {
76        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
77            // Try L1EvictUnchanged
78            {
79                let saved_pos = stream.position();
80                if stream.expect_string(".L1::evict_unchanged").is_ok() {
81                    return Ok(Level1EvictionPriority::L1EvictUnchanged);
82                }
83                stream.set_position(saved_pos);
84            }
85            let saved_pos = stream.position();
86            // Try L1EvictNormal
87            {
88                let saved_pos = stream.position();
89                if stream.expect_string(".L1::evict_normal").is_ok() {
90                    return Ok(Level1EvictionPriority::L1EvictNormal);
91                }
92                stream.set_position(saved_pos);
93            }
94            stream.set_position(saved_pos);
95            let saved_pos = stream.position();
96            // Try L1EvictFirst
97            {
98                let saved_pos = stream.position();
99                if stream.expect_string(".L1::evict_first").is_ok() {
100                    return Ok(Level1EvictionPriority::L1EvictFirst);
101                }
102                stream.set_position(saved_pos);
103            }
104            stream.set_position(saved_pos);
105            let saved_pos = stream.position();
106            // Try L1NoAllocate
107            {
108                let saved_pos = stream.position();
109                if stream.expect_string(".L1::no_allocate").is_ok() {
110                    return Ok(Level1EvictionPriority::L1NoAllocate);
111                }
112                stream.set_position(saved_pos);
113            }
114            stream.set_position(saved_pos);
115            let saved_pos = stream.position();
116            // Try L1EvictLast
117            {
118                let saved_pos = stream.position();
119                if stream.expect_string(".L1::evict_last").is_ok() {
120                    return Ok(Level1EvictionPriority::L1EvictLast);
121                }
122                stream.set_position(saved_pos);
123            }
124            stream.set_position(saved_pos);
125            let span = stream
126                .peek()
127                .map(|(_, s)| s.clone())
128                .unwrap_or(Span { start: 0, end: 0 });
129            let expected = &[
130                ".L1::evict_unchanged",
131                ".L1::evict_normal",
132                ".L1::evict_first",
133                ".L1::no_allocate",
134                ".L1::evict_last",
135            ];
136            let found = stream
137                .peek()
138                .map(|(t, _)| format!("{:?}", t))
139                .unwrap_or_else(|_| "<end of input>".to_string());
140            Err(crate::parser::unexpected_value(span, expected, found))
141        }
142    }
143
144    impl PtxParser for Level2EvictionPriority {
145        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
146            // Try L2EvictNormal
147            {
148                let saved_pos = stream.position();
149                if stream.expect_string(".L2::evict_normal").is_ok() {
150                    return Ok(Level2EvictionPriority::L2EvictNormal);
151                }
152                stream.set_position(saved_pos);
153            }
154            let saved_pos = stream.position();
155            // Try L2EvictFirst
156            {
157                let saved_pos = stream.position();
158                if stream.expect_string(".L2::evict_first").is_ok() {
159                    return Ok(Level2EvictionPriority::L2EvictFirst);
160                }
161                stream.set_position(saved_pos);
162            }
163            stream.set_position(saved_pos);
164            let saved_pos = stream.position();
165            // Try L2EvictLast
166            {
167                let saved_pos = stream.position();
168                if stream.expect_string(".L2::evict_last").is_ok() {
169                    return Ok(Level2EvictionPriority::L2EvictLast);
170                }
171                stream.set_position(saved_pos);
172            }
173            stream.set_position(saved_pos);
174            let span = stream
175                .peek()
176                .map(|(_, s)| s.clone())
177                .unwrap_or(Span { start: 0, end: 0 });
178            let expected = &[".L2::evict_normal", ".L2::evict_first", ".L2::evict_last"];
179            let found = stream
180                .peek()
181                .map(|(t, _)| format!("{:?}", t))
182                .unwrap_or_else(|_| "<end of input>".to_string());
183            Err(crate::parser::unexpected_value(span, expected, found))
184        }
185    }
186
187    impl PtxParser for LevelCacheHint {
188        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
189            // Try L2CacheHint
190            {
191                let saved_pos = stream.position();
192                if stream.expect_string(".L2::cache_hint").is_ok() {
193                    return Ok(LevelCacheHint::L2CacheHint);
194                }
195                stream.set_position(saved_pos);
196            }
197            let span = stream
198                .peek()
199                .map(|(_, s)| s.clone())
200                .unwrap_or(Span { start: 0, end: 0 });
201            let expected = &[".L2::cache_hint"];
202            let found = stream
203                .peek()
204                .map(|(t, _)| format!("{:?}", t))
205                .unwrap_or_else(|_| "<end of input>".to_string());
206            Err(crate::parser::unexpected_value(span, expected, found))
207        }
208    }
209
210    impl PtxParser for LevelPrefetchSize {
211        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
212            // Try L2128b
213            {
214                let saved_pos = stream.position();
215                if stream.expect_string(".L2::128B").is_ok() {
216                    return Ok(LevelPrefetchSize::L2128b);
217                }
218                stream.set_position(saved_pos);
219            }
220            let saved_pos = stream.position();
221            // Try L2256b
222            {
223                let saved_pos = stream.position();
224                if stream.expect_string(".L2::256B").is_ok() {
225                    return Ok(LevelPrefetchSize::L2256b);
226                }
227                stream.set_position(saved_pos);
228            }
229            stream.set_position(saved_pos);
230            let saved_pos = stream.position();
231            // Try L264b
232            {
233                let saved_pos = stream.position();
234                if stream.expect_string(".L2::64B").is_ok() {
235                    return Ok(LevelPrefetchSize::L264b);
236                }
237                stream.set_position(saved_pos);
238            }
239            stream.set_position(saved_pos);
240            let span = stream
241                .peek()
242                .map(|(_, s)| s.clone())
243                .unwrap_or(Span { start: 0, end: 0 });
244            let expected = &[".L2::128B", ".L2::256B", ".L2::64B"];
245            let found = stream
246                .peek()
247                .map(|(t, _)| format!("{:?}", t))
248                .unwrap_or_else(|_| "<end of input>".to_string());
249            Err(crate::parser::unexpected_value(span, expected, found))
250        }
251    }
252
253    impl PtxParser for Type {
254        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
255            // Try B128
256            {
257                let saved_pos = stream.position();
258                if stream.expect_string(".b128").is_ok() {
259                    return Ok(Type::B128);
260                }
261                stream.set_position(saved_pos);
262            }
263            let saved_pos = stream.position();
264            // Try B16
265            {
266                let saved_pos = stream.position();
267                if stream.expect_string(".b16").is_ok() {
268                    return Ok(Type::B16);
269                }
270                stream.set_position(saved_pos);
271            }
272            stream.set_position(saved_pos);
273            let saved_pos = stream.position();
274            // Try B32
275            {
276                let saved_pos = stream.position();
277                if stream.expect_string(".b32").is_ok() {
278                    return Ok(Type::B32);
279                }
280                stream.set_position(saved_pos);
281            }
282            stream.set_position(saved_pos);
283            let saved_pos = stream.position();
284            // Try B64
285            {
286                let saved_pos = stream.position();
287                if stream.expect_string(".b64").is_ok() {
288                    return Ok(Type::B64);
289                }
290                stream.set_position(saved_pos);
291            }
292            stream.set_position(saved_pos);
293            let saved_pos = stream.position();
294            // Try U16
295            {
296                let saved_pos = stream.position();
297                if stream.expect_string(".u16").is_ok() {
298                    return Ok(Type::U16);
299                }
300                stream.set_position(saved_pos);
301            }
302            stream.set_position(saved_pos);
303            let saved_pos = stream.position();
304            // Try U32
305            {
306                let saved_pos = stream.position();
307                if stream.expect_string(".u32").is_ok() {
308                    return Ok(Type::U32);
309                }
310                stream.set_position(saved_pos);
311            }
312            stream.set_position(saved_pos);
313            let saved_pos = stream.position();
314            // Try U64
315            {
316                let saved_pos = stream.position();
317                if stream.expect_string(".u64").is_ok() {
318                    return Ok(Type::U64);
319                }
320                stream.set_position(saved_pos);
321            }
322            stream.set_position(saved_pos);
323            let saved_pos = stream.position();
324            // Try S16
325            {
326                let saved_pos = stream.position();
327                if stream.expect_string(".s16").is_ok() {
328                    return Ok(Type::S16);
329                }
330                stream.set_position(saved_pos);
331            }
332            stream.set_position(saved_pos);
333            let saved_pos = stream.position();
334            // Try S32
335            {
336                let saved_pos = stream.position();
337                if stream.expect_string(".s32").is_ok() {
338                    return Ok(Type::S32);
339                }
340                stream.set_position(saved_pos);
341            }
342            stream.set_position(saved_pos);
343            let saved_pos = stream.position();
344            // Try S64
345            {
346                let saved_pos = stream.position();
347                if stream.expect_string(".s64").is_ok() {
348                    return Ok(Type::S64);
349                }
350                stream.set_position(saved_pos);
351            }
352            stream.set_position(saved_pos);
353            let saved_pos = stream.position();
354            // Try F32
355            {
356                let saved_pos = stream.position();
357                if stream.expect_string(".f32").is_ok() {
358                    return Ok(Type::F32);
359                }
360                stream.set_position(saved_pos);
361            }
362            stream.set_position(saved_pos);
363            let saved_pos = stream.position();
364            // Try F64
365            {
366                let saved_pos = stream.position();
367                if stream.expect_string(".f64").is_ok() {
368                    return Ok(Type::F64);
369                }
370                stream.set_position(saved_pos);
371            }
372            stream.set_position(saved_pos);
373            let saved_pos = stream.position();
374            // Try B8
375            {
376                let saved_pos = stream.position();
377                if stream.expect_string(".b8").is_ok() {
378                    return Ok(Type::B8);
379                }
380                stream.set_position(saved_pos);
381            }
382            stream.set_position(saved_pos);
383            let saved_pos = stream.position();
384            // Try U8
385            {
386                let saved_pos = stream.position();
387                if stream.expect_string(".u8").is_ok() {
388                    return Ok(Type::U8);
389                }
390                stream.set_position(saved_pos);
391            }
392            stream.set_position(saved_pos);
393            let saved_pos = stream.position();
394            // Try S8
395            {
396                let saved_pos = stream.position();
397                if stream.expect_string(".s8").is_ok() {
398                    return Ok(Type::S8);
399                }
400                stream.set_position(saved_pos);
401            }
402            stream.set_position(saved_pos);
403            let span = stream
404                .peek()
405                .map(|(_, s)| s.clone())
406                .unwrap_or(Span { start: 0, end: 0 });
407            let expected = &[
408                ".b128", ".b16", ".b32", ".b64", ".u16", ".u32", ".u64", ".s16", ".s32", ".s64",
409                ".f32", ".f64", ".b8", ".u8", ".s8",
410            ];
411            let found = stream
412                .peek()
413                .map(|(t, _)| format!("{:?}", t))
414                .unwrap_or_else(|_| "<end of input>".to_string());
415            Err(crate::parser::unexpected_value(span, expected, found))
416        }
417    }
418
419    impl PtxParser for Vec {
420        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
421            // Try V2
422            {
423                let saved_pos = stream.position();
424                if stream.expect_string(".v2").is_ok() {
425                    return Ok(Vec::V2);
426                }
427                stream.set_position(saved_pos);
428            }
429            let saved_pos = stream.position();
430            // Try V4
431            {
432                let saved_pos = stream.position();
433                if stream.expect_string(".v4").is_ok() {
434                    return Ok(Vec::V4);
435                }
436                stream.set_position(saved_pos);
437            }
438            stream.set_position(saved_pos);
439            let saved_pos = stream.position();
440            // Try V8
441            {
442                let saved_pos = stream.position();
443                if stream.expect_string(".v8").is_ok() {
444                    return Ok(Vec::V8);
445                }
446                stream.set_position(saved_pos);
447            }
448            stream.set_position(saved_pos);
449            let span = stream
450                .peek()
451                .map(|(_, s)| s.clone())
452                .unwrap_or(Span { start: 0, end: 0 });
453            let expected = &[".v2", ".v4", ".v8"];
454            let found = stream
455                .peek()
456                .map(|(t, _)| format!("{:?}", t))
457                .unwrap_or_else(|_| "<end of input>".to_string());
458            Err(crate::parser::unexpected_value(span, expected, found))
459        }
460    }
461
462    impl PtxParser for LdGlobalCopNcLevelCacheHintLevelPrefetchSizeType {
463        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
464            stream.expect_string("ld")?;
465            stream.expect_string(".global")?;
466            let global = ();
467            stream.expect_complete()?;
468            let saved_pos = stream.position();
469            let cop = match Cop::parse(stream) {
470                Ok(val) => Some(val),
471                Err(_) => {
472                    stream.set_position(saved_pos);
473                    None
474                }
475            };
476            stream.expect_complete()?;
477            stream.expect_string(".nc")?;
478            let nc = ();
479            stream.expect_complete()?;
480            let saved_pos = stream.position();
481            let level_cache_hint = match LevelCacheHint::parse(stream) {
482                Ok(val) => Some(val),
483                Err(_) => {
484                    stream.set_position(saved_pos);
485                    None
486                }
487            };
488            stream.expect_complete()?;
489            let saved_pos = stream.position();
490            let level_prefetch_size = match LevelPrefetchSize::parse(stream) {
491                Ok(val) => Some(val),
492                Err(_) => {
493                    stream.set_position(saved_pos);
494                    None
495                }
496            };
497            stream.expect_complete()?;
498            let type_ = Type::parse(stream)?;
499            stream.expect_complete()?;
500            let d = GeneralOperand::parse(stream)?;
501            stream.expect_complete()?;
502            stream.expect(&PtxToken::Comma)?;
503            let a = AddressOperand::parse(stream)?;
504            stream.expect_complete()?;
505            let saved_pos = stream.position();
506            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
507            if !has_comma {
508                stream.set_position(saved_pos);
509            }
510            let saved_pos = stream.position();
511            let cache_policy = match GeneralOperand::parse(stream) {
512                Ok(val) => Some(val),
513                Err(_) => {
514                    stream.set_position(saved_pos);
515                    None
516                }
517            };
518            stream.expect_complete()?;
519            stream.expect_complete()?;
520            stream.expect(&PtxToken::Semicolon)?;
521            Ok(LdGlobalCopNcLevelCacheHintLevelPrefetchSizeType {
522                global,
523                cop,
524                nc,
525                level_cache_hint,
526                level_prefetch_size,
527                type_,
528                d,
529                a,
530                cache_policy,
531            })
532        }
533    }
534
535    impl PtxParser for LdGlobalCopNcLevelCacheHintLevelPrefetchSizeVecType {
536        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
537            stream.expect_string("ld")?;
538            stream.expect_string(".global")?;
539            let global = ();
540            stream.expect_complete()?;
541            let saved_pos = stream.position();
542            let cop = match Cop::parse(stream) {
543                Ok(val) => Some(val),
544                Err(_) => {
545                    stream.set_position(saved_pos);
546                    None
547                }
548            };
549            stream.expect_complete()?;
550            stream.expect_string(".nc")?;
551            let nc = ();
552            stream.expect_complete()?;
553            let saved_pos = stream.position();
554            let level_cache_hint = match LevelCacheHint::parse(stream) {
555                Ok(val) => Some(val),
556                Err(_) => {
557                    stream.set_position(saved_pos);
558                    None
559                }
560            };
561            stream.expect_complete()?;
562            let saved_pos = stream.position();
563            let level_prefetch_size = match LevelPrefetchSize::parse(stream) {
564                Ok(val) => Some(val),
565                Err(_) => {
566                    stream.set_position(saved_pos);
567                    None
568                }
569            };
570            stream.expect_complete()?;
571            let vec = Vec::parse(stream)?;
572            stream.expect_complete()?;
573            let type_ = Type::parse(stream)?;
574            stream.expect_complete()?;
575            let d = GeneralOperand::parse(stream)?;
576            stream.expect_complete()?;
577            stream.expect(&PtxToken::Comma)?;
578            let a = AddressOperand::parse(stream)?;
579            stream.expect_complete()?;
580            let saved_pos = stream.position();
581            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
582            if !has_comma {
583                stream.set_position(saved_pos);
584            }
585            let saved_pos = stream.position();
586            let cache_policy = match GeneralOperand::parse(stream) {
587                Ok(val) => Some(val),
588                Err(_) => {
589                    stream.set_position(saved_pos);
590                    None
591                }
592            };
593            stream.expect_complete()?;
594            stream.expect_complete()?;
595            stream.expect(&PtxToken::Semicolon)?;
596            Ok(LdGlobalCopNcLevelCacheHintLevelPrefetchSizeVecType {
597                global,
598                cop,
599                nc,
600                level_cache_hint,
601                level_prefetch_size,
602                vec,
603                type_,
604                d,
605                a,
606                cache_policy,
607            })
608        }
609    }
610
611    impl PtxParser for LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeType {
612        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
613            stream.expect_string("ld")?;
614            stream.expect_string(".global")?;
615            let global = ();
616            stream.expect_complete()?;
617            stream.expect_string(".nc")?;
618            let nc = ();
619            stream.expect_complete()?;
620            let saved_pos = stream.position();
621            let level1_eviction_priority = match Level1EvictionPriority::parse(stream) {
622                Ok(val) => Some(val),
623                Err(_) => {
624                    stream.set_position(saved_pos);
625                    None
626                }
627            };
628            stream.expect_complete()?;
629            let saved_pos = stream.position();
630            let level2_eviction_priority = match Level2EvictionPriority::parse(stream) {
631                Ok(val) => Some(val),
632                Err(_) => {
633                    stream.set_position(saved_pos);
634                    None
635                }
636            };
637            stream.expect_complete()?;
638            let saved_pos = stream.position();
639            let level_cache_hint = match LevelCacheHint::parse(stream) {
640                Ok(val) => Some(val),
641                Err(_) => {
642                    stream.set_position(saved_pos);
643                    None
644                }
645            };
646            stream.expect_complete()?;
647            let saved_pos = stream.position();
648            let level_prefetch_size = match LevelPrefetchSize::parse(stream) {
649                Ok(val) => Some(val),
650                Err(_) => {
651                    stream.set_position(saved_pos);
652                    None
653                }
654            };
655            stream.expect_complete()?;
656            let type_ = Type::parse(stream)?;
657            stream.expect_complete()?;
658            let d = GeneralOperand::parse(stream)?;
659            stream.expect_complete()?;
660            stream.expect(&PtxToken::Comma)?;
661            let a = AddressOperand::parse(stream)?;
662            stream.expect_complete()?;
663            let saved_pos = stream.position();
664            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
665            if !has_comma {
666                stream.set_position(saved_pos);
667            }
668            let saved_pos = stream.position();
669            let cache_policy = match GeneralOperand::parse(stream) {
670                Ok(val) => Some(val),
671                Err(_) => {
672                    stream.set_position(saved_pos);
673                    None
674                }
675            };
676            stream.expect_complete()?;
677            stream.expect_complete()?;
678            stream.expect(&PtxToken::Semicolon)?;
679            Ok(LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeType {
680                global,
681                nc,
682                level1_eviction_priority,
683                level2_eviction_priority,
684                level_cache_hint,
685                level_prefetch_size,
686                type_,
687                d,
688                a,
689                cache_policy,
690            })
691        }
692    }
693
694    impl PtxParser for LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType {
695        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
696            stream.expect_string("ld")?;
697            stream.expect_string(".global")?;
698            let global = ();
699            stream.expect_complete()?;
700            stream.expect_string(".nc")?;
701            let nc = ();
702            stream.expect_complete()?;
703            let saved_pos = stream.position();
704            let level1_eviction_priority = match Level1EvictionPriority::parse(stream) {
705                Ok(val) => Some(val),
706                Err(_) => {
707                    stream.set_position(saved_pos);
708                    None
709                }
710            };
711            stream.expect_complete()?;
712            let saved_pos = stream.position();
713            let level2_eviction_priority = match Level2EvictionPriority::parse(stream) {
714                Ok(val) => Some(val),
715                Err(_) => {
716                    stream.set_position(saved_pos);
717                    None
718                }
719            };
720            stream.expect_complete()?;
721            let saved_pos = stream.position();
722            let level_cache_hint = match LevelCacheHint::parse(stream) {
723                Ok(val) => Some(val),
724                Err(_) => {
725                    stream.set_position(saved_pos);
726                    None
727                }
728            };
729            stream.expect_complete()?;
730            let saved_pos = stream.position();
731            let level_prefetch_size = match LevelPrefetchSize::parse(stream) {
732                Ok(val) => Some(val),
733                Err(_) => {
734                    stream.set_position(saved_pos);
735                    None
736                }
737            };
738            stream.expect_complete()?;
739            let vec = Vec::parse(stream)?;
740            stream.expect_complete()?;
741            let type_ = Type::parse(stream)?;
742            stream.expect_complete()?;
743            let d = GeneralOperand::parse(stream)?;
744            stream.expect_complete()?;
745            stream.expect(&PtxToken::Comma)?;
746            let a = AddressOperand::parse(stream)?;
747            stream.expect_complete()?;
748            let saved_pos = stream.position();
749            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
750            if !has_comma {
751                stream.set_position(saved_pos);
752            }
753            let saved_pos = stream.position();
754            let cache_policy = match GeneralOperand::parse(stream) {
755                Ok(val) => Some(val),
756                Err(_) => {
757                    stream.set_position(saved_pos);
758                    None
759                }
760            };
761            stream.expect_complete()?;
762            stream.expect_complete()?;
763            stream.expect(&PtxToken::Semicolon)?;
764            Ok(LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType {
765                global,
766                nc,
767                level1_eviction_priority,
768                level2_eviction_priority,
769                level_cache_hint,
770                level_prefetch_size,
771                vec,
772                type_,
773                d,
774                a,
775                cache_policy,
776            })
777        }
778    }
779}
ptx_parser/parser/instruction/ld_global_nc.rs

ptx_parser/parser/instruction/
ld_global_nc.rs
