ptx_parser/parser/instruction/
cp_async_bulk.rs

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

ptx_parser/parser/instruction/
cp_async_bulk.rs
