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

ptx_parser/parser/instruction/
cp_async_bulk.rs
