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

ptx_parser/parser/instruction/
ld_global_nc.rs
