ptx_parser/parser/instruction/
cp_async_bulk_tensor.rs

1//! Original PTX specification:
2//!
3//! // global -> shared::cta
4//! cp.async.bulk.tensor.dim.dst.src{.load_mode}.completion_mechanism{.cta_group}{.level::cache_hint} [dstMem], [tensorMap, tensorCoords], [mbar]{, im2colInfo} {, cache-policy};
5//! .dst =                  { .shared::cta };
6//! .src =                  { .global };
7//! .dim =                  { .1d, .2d, .3d, .4d, .5d };
8//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
9//! .cta_group =            { .cta_group::1, .cta_group::2 };
10//! .load_mode =            { .tile, .tile::gather4, .im2col, .im2col::w, .im2col::w::128 };
11//! .level::cache_hint =    { .L2::cache_hint };
12//! ----------------------------------------------------------------
13//! // global -> shared::cluster
14//! cp.async.bulk.tensor.dim.dst.src{.load_mode}.completion_mechanism{.multicast}{.cta_group}{.level::cache_hint} [dstMem], [tensorMap, tensorCoords], [mbar]{, im2colInfo} {, ctaMask} {, cache-policy};
15//! .dst =                  { .shared::cluster };
16//! .src =                  { .global };
17//! .dim =                  { .1d, .2d, .3d, .4d, .5d };
18//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
19//! .cta_group =            { .cta_group::1, .cta_group::2 };
20//! .load_mode =            { .tile, .tile::gather4, .im2col, .im2col::w, .im2col::w::128 };
21//! .level::cache_hint =    { .L2::cache_hint };
22//! .multicast =            { .multicast::cluster  };
23//! ----------------------------------------------------------------
24//! // shared::cta -> global;
25//! cp.async.bulk.tensor.dim.dst.src{.load_mode}.completion_mechanism{.level::cache_hint} [tensorMap, tensorCoords], [srcMem] {, cache-policy};
26//! .dst =                  { .global };
27//! .src =                  { .shared::cta };
28//! .dim =                  { .1d, .2d, .3d, .4d, .5d };
29//! .completion_mechanism = { .bulk_group };
30//! .load_mode =            { .tile, .tile::scatter4, .im2col_no_offs };
31//! .level::cache_hint =    { .L2::cache_hint };
32
33#![allow(unused)]
34
35use crate::lexer::PtxToken;
36use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
37use crate::r#type::common::*;
38
39pub mod section_0 {
40    use super::*;
41    use crate::r#type::instruction::cp_async_bulk_tensor::section_0::*;
42
43    // ============================================================================
44    // Generated enum parsers
45    // ============================================================================
46
47    impl PtxParser for CompletionMechanism {
48        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
49            // Try MbarrierCompleteTxBytes
50            {
51                let saved_pos = stream.position();
52                if stream.expect_string(".mbarrier::complete_tx::bytes").is_ok() {
53                    return Ok(CompletionMechanism::MbarrierCompleteTxBytes);
54                }
55                stream.set_position(saved_pos);
56            }
57            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
58            let expected = &[".mbarrier::complete_tx::bytes"];
59            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
60            Err(crate::parser::unexpected_value(span, expected, found))
61        }
62    }
63
64    impl PtxParser for CtaGroup {
65        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
66            // Try CtaGroup1
67            {
68                let saved_pos = stream.position();
69                if stream.expect_string(".cta_group::1").is_ok() {
70                    return Ok(CtaGroup::CtaGroup1);
71                }
72                stream.set_position(saved_pos);
73            }
74            let saved_pos = stream.position();
75            // Try CtaGroup2
76            {
77                let saved_pos = stream.position();
78                if stream.expect_string(".cta_group::2").is_ok() {
79                    return Ok(CtaGroup::CtaGroup2);
80                }
81                stream.set_position(saved_pos);
82            }
83            stream.set_position(saved_pos);
84            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
85            let expected = &[".cta_group::1", ".cta_group::2"];
86            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
87            Err(crate::parser::unexpected_value(span, expected, found))
88        }
89    }
90
91    impl PtxParser for Dim {
92        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
93            // Try _1d
94            {
95                let saved_pos = stream.position();
96                if stream.expect_string(".1d").is_ok() {
97                    return Ok(Dim::_1d);
98                }
99                stream.set_position(saved_pos);
100            }
101            let saved_pos = stream.position();
102            // Try _2d
103            {
104                let saved_pos = stream.position();
105                if stream.expect_string(".2d").is_ok() {
106                    return Ok(Dim::_2d);
107                }
108                stream.set_position(saved_pos);
109            }
110            stream.set_position(saved_pos);
111            let saved_pos = stream.position();
112            // Try _3d
113            {
114                let saved_pos = stream.position();
115                if stream.expect_string(".3d").is_ok() {
116                    return Ok(Dim::_3d);
117                }
118                stream.set_position(saved_pos);
119            }
120            stream.set_position(saved_pos);
121            let saved_pos = stream.position();
122            // Try _4d
123            {
124                let saved_pos = stream.position();
125                if stream.expect_string(".4d").is_ok() {
126                    return Ok(Dim::_4d);
127                }
128                stream.set_position(saved_pos);
129            }
130            stream.set_position(saved_pos);
131            let saved_pos = stream.position();
132            // Try _5d
133            {
134                let saved_pos = stream.position();
135                if stream.expect_string(".5d").is_ok() {
136                    return Ok(Dim::_5d);
137                }
138                stream.set_position(saved_pos);
139            }
140            stream.set_position(saved_pos);
141            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
142            let expected = &[".1d", ".2d", ".3d", ".4d", ".5d"];
143            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
144            Err(crate::parser::unexpected_value(span, expected, found))
145        }
146    }
147
148    impl PtxParser for Dst {
149        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
150            // Try SharedCta
151            {
152                let saved_pos = stream.position();
153                if stream.expect_string(".shared::cta").is_ok() {
154                    return Ok(Dst::SharedCta);
155                }
156                stream.set_position(saved_pos);
157            }
158            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
159            let expected = &[".shared::cta"];
160            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
161            Err(crate::parser::unexpected_value(span, expected, found))
162        }
163    }
164
165    impl PtxParser for LevelCacheHint {
166        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
167            // Try L2CacheHint
168            {
169                let saved_pos = stream.position();
170                if stream.expect_string(".L2::cache_hint").is_ok() {
171                    return Ok(LevelCacheHint::L2CacheHint);
172                }
173                stream.set_position(saved_pos);
174            }
175            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
176            let expected = &[".L2::cache_hint"];
177            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
178            Err(crate::parser::unexpected_value(span, expected, found))
179        }
180    }
181
182    impl PtxParser for LoadMode {
183        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
184            // Try Im2colW128
185            {
186                let saved_pos = stream.position();
187                if stream.expect_string(".im2col::w::128").is_ok() {
188                    return Ok(LoadMode::Im2colW128);
189                }
190                stream.set_position(saved_pos);
191            }
192            let saved_pos = stream.position();
193            // Try TileGather4
194            {
195                let saved_pos = stream.position();
196                if stream.expect_string(".tile::gather4").is_ok() {
197                    return Ok(LoadMode::TileGather4);
198                }
199                stream.set_position(saved_pos);
200            }
201            stream.set_position(saved_pos);
202            let saved_pos = stream.position();
203            // Try Im2colW
204            {
205                let saved_pos = stream.position();
206                if stream.expect_string(".im2col::w").is_ok() {
207                    return Ok(LoadMode::Im2colW);
208                }
209                stream.set_position(saved_pos);
210            }
211            stream.set_position(saved_pos);
212            let saved_pos = stream.position();
213            // Try Im2col
214            {
215                let saved_pos = stream.position();
216                if stream.expect_string(".im2col").is_ok() {
217                    return Ok(LoadMode::Im2col);
218                }
219                stream.set_position(saved_pos);
220            }
221            stream.set_position(saved_pos);
222            let saved_pos = stream.position();
223            // Try Tile
224            {
225                let saved_pos = stream.position();
226                if stream.expect_string(".tile").is_ok() {
227                    return Ok(LoadMode::Tile);
228                }
229                stream.set_position(saved_pos);
230            }
231            stream.set_position(saved_pos);
232            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
233            let expected = &[".im2col::w::128", ".tile::gather4", ".im2col::w", ".im2col", ".tile"];
234            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
235            Err(crate::parser::unexpected_value(span, expected, found))
236        }
237    }
238
239    impl PtxParser for Src {
240        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
241            // Try Global
242            {
243                let saved_pos = stream.position();
244                if stream.expect_string(".global").is_ok() {
245                    return Ok(Src::Global);
246                }
247                stream.set_position(saved_pos);
248            }
249            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
250            let expected = &[".global"];
251            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
252            Err(crate::parser::unexpected_value(span, expected, found))
253        }
254    }
255
256    impl PtxParser for CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismCtaGroupLevelCacheHint {
257        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
258            stream.expect_string("cp")?;
259            stream.expect_string(".async")?;
260            let async_ = ();
261            stream.expect_complete()?;
262            stream.expect_string(".bulk")?;
263            let bulk = ();
264            stream.expect_complete()?;
265            stream.expect_string(".tensor")?;
266            let tensor = ();
267            stream.expect_complete()?;
268            let dim = Dim::parse(stream)?;
269            stream.expect_complete()?;
270            let dst = Dst::parse(stream)?;
271            stream.expect_complete()?;
272            let src = Src::parse(stream)?;
273            stream.expect_complete()?;
274            let saved_pos = stream.position();
275            let load_mode = match LoadMode::parse(stream) {
276                Ok(val) => Some(val),
277                Err(_) => {
278                    stream.set_position(saved_pos);
279                    None
280                }
281            };
282            stream.expect_complete()?;
283            let completion_mechanism = CompletionMechanism::parse(stream)?;
284            stream.expect_complete()?;
285            let saved_pos = stream.position();
286            let cta_group = match CtaGroup::parse(stream) {
287                Ok(val) => Some(val),
288                Err(_) => {
289                    stream.set_position(saved_pos);
290                    None
291                }
292            };
293            stream.expect_complete()?;
294            let saved_pos = stream.position();
295            let level_cache_hint = match LevelCacheHint::parse(stream) {
296                Ok(val) => Some(val),
297                Err(_) => {
298                    stream.set_position(saved_pos);
299                    None
300                }
301            };
302            stream.expect_complete()?;
303            let dstmem = AddressOperand::parse(stream)?;
304            stream.expect_complete()?;
305            stream.expect(&PtxToken::Comma)?;
306            let tensormap = TexHandler2::parse(stream)?;
307            stream.expect_complete()?;
308            stream.expect(&PtxToken::Comma)?;
309            let mbar = AddressOperand::parse(stream)?;
310            stream.expect_complete()?;
311            let saved_pos = stream.position();
312            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
313            if !has_comma {
314                stream.set_position(saved_pos);
315            }
316            let saved_pos = stream.position();
317            let im2colinfo = match GeneralOperand::parse(stream) {
318                Ok(val) => Some(val),
319                Err(_) => {
320                    stream.set_position(saved_pos);
321                    None
322                }
323            };
324            stream.expect_complete()?;
325            let saved_pos = stream.position();
326            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
327            if !has_comma {
328                stream.set_position(saved_pos);
329            }
330            let saved_pos = stream.position();
331            let cache_policy = match GeneralOperand::parse(stream) {
332                Ok(val) => Some(val),
333                Err(_) => {
334                    stream.set_position(saved_pos);
335                    None
336                }
337            };
338            stream.expect_complete()?;
339            stream.expect_complete()?;
340            stream.expect(&PtxToken::Semicolon)?;
341            Ok(CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismCtaGroupLevelCacheHint {
342                async_,
343                bulk,
344                tensor,
345                dim,
346                dst,
347                src,
348                load_mode,
349                completion_mechanism,
350                cta_group,
351                level_cache_hint,
352                dstmem,
353                tensormap,
354                mbar,
355                im2colinfo,
356                cache_policy,
357            })
358        }
359    }
360
361
362}
363
364pub mod section_1 {
365    use super::*;
366    use crate::r#type::instruction::cp_async_bulk_tensor::section_1::*;
367
368    // ============================================================================
369    // Generated enum parsers
370    // ============================================================================
371
372    impl PtxParser for CompletionMechanism {
373        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
374            // Try MbarrierCompleteTxBytes
375            {
376                let saved_pos = stream.position();
377                if stream.expect_string(".mbarrier::complete_tx::bytes").is_ok() {
378                    return Ok(CompletionMechanism::MbarrierCompleteTxBytes);
379                }
380                stream.set_position(saved_pos);
381            }
382            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
383            let expected = &[".mbarrier::complete_tx::bytes"];
384            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
385            Err(crate::parser::unexpected_value(span, expected, found))
386        }
387    }
388
389    impl PtxParser for CtaGroup {
390        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
391            // Try CtaGroup1
392            {
393                let saved_pos = stream.position();
394                if stream.expect_string(".cta_group::1").is_ok() {
395                    return Ok(CtaGroup::CtaGroup1);
396                }
397                stream.set_position(saved_pos);
398            }
399            let saved_pos = stream.position();
400            // Try CtaGroup2
401            {
402                let saved_pos = stream.position();
403                if stream.expect_string(".cta_group::2").is_ok() {
404                    return Ok(CtaGroup::CtaGroup2);
405                }
406                stream.set_position(saved_pos);
407            }
408            stream.set_position(saved_pos);
409            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
410            let expected = &[".cta_group::1", ".cta_group::2"];
411            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
412            Err(crate::parser::unexpected_value(span, expected, found))
413        }
414    }
415
416    impl PtxParser for Dim {
417        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
418            // Try _1d
419            {
420                let saved_pos = stream.position();
421                if stream.expect_string(".1d").is_ok() {
422                    return Ok(Dim::_1d);
423                }
424                stream.set_position(saved_pos);
425            }
426            let saved_pos = stream.position();
427            // Try _2d
428            {
429                let saved_pos = stream.position();
430                if stream.expect_string(".2d").is_ok() {
431                    return Ok(Dim::_2d);
432                }
433                stream.set_position(saved_pos);
434            }
435            stream.set_position(saved_pos);
436            let saved_pos = stream.position();
437            // Try _3d
438            {
439                let saved_pos = stream.position();
440                if stream.expect_string(".3d").is_ok() {
441                    return Ok(Dim::_3d);
442                }
443                stream.set_position(saved_pos);
444            }
445            stream.set_position(saved_pos);
446            let saved_pos = stream.position();
447            // Try _4d
448            {
449                let saved_pos = stream.position();
450                if stream.expect_string(".4d").is_ok() {
451                    return Ok(Dim::_4d);
452                }
453                stream.set_position(saved_pos);
454            }
455            stream.set_position(saved_pos);
456            let saved_pos = stream.position();
457            // Try _5d
458            {
459                let saved_pos = stream.position();
460                if stream.expect_string(".5d").is_ok() {
461                    return Ok(Dim::_5d);
462                }
463                stream.set_position(saved_pos);
464            }
465            stream.set_position(saved_pos);
466            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
467            let expected = &[".1d", ".2d", ".3d", ".4d", ".5d"];
468            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
469            Err(crate::parser::unexpected_value(span, expected, found))
470        }
471    }
472
473    impl PtxParser for Dst {
474        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
475            // Try SharedCluster
476            {
477                let saved_pos = stream.position();
478                if stream.expect_string(".shared::cluster").is_ok() {
479                    return Ok(Dst::SharedCluster);
480                }
481                stream.set_position(saved_pos);
482            }
483            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
484            let expected = &[".shared::cluster"];
485            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
486            Err(crate::parser::unexpected_value(span, expected, found))
487        }
488    }
489
490    impl PtxParser for LevelCacheHint {
491        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
492            // Try L2CacheHint
493            {
494                let saved_pos = stream.position();
495                if stream.expect_string(".L2::cache_hint").is_ok() {
496                    return Ok(LevelCacheHint::L2CacheHint);
497                }
498                stream.set_position(saved_pos);
499            }
500            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
501            let expected = &[".L2::cache_hint"];
502            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
503            Err(crate::parser::unexpected_value(span, expected, found))
504        }
505    }
506
507    impl PtxParser for LoadMode {
508        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
509            // Try Im2colW128
510            {
511                let saved_pos = stream.position();
512                if stream.expect_string(".im2col::w::128").is_ok() {
513                    return Ok(LoadMode::Im2colW128);
514                }
515                stream.set_position(saved_pos);
516            }
517            let saved_pos = stream.position();
518            // Try TileGather4
519            {
520                let saved_pos = stream.position();
521                if stream.expect_string(".tile::gather4").is_ok() {
522                    return Ok(LoadMode::TileGather4);
523                }
524                stream.set_position(saved_pos);
525            }
526            stream.set_position(saved_pos);
527            let saved_pos = stream.position();
528            // Try Im2colW
529            {
530                let saved_pos = stream.position();
531                if stream.expect_string(".im2col::w").is_ok() {
532                    return Ok(LoadMode::Im2colW);
533                }
534                stream.set_position(saved_pos);
535            }
536            stream.set_position(saved_pos);
537            let saved_pos = stream.position();
538            // Try Im2col
539            {
540                let saved_pos = stream.position();
541                if stream.expect_string(".im2col").is_ok() {
542                    return Ok(LoadMode::Im2col);
543                }
544                stream.set_position(saved_pos);
545            }
546            stream.set_position(saved_pos);
547            let saved_pos = stream.position();
548            // Try Tile
549            {
550                let saved_pos = stream.position();
551                if stream.expect_string(".tile").is_ok() {
552                    return Ok(LoadMode::Tile);
553                }
554                stream.set_position(saved_pos);
555            }
556            stream.set_position(saved_pos);
557            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
558            let expected = &[".im2col::w::128", ".tile::gather4", ".im2col::w", ".im2col", ".tile"];
559            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
560            Err(crate::parser::unexpected_value(span, expected, found))
561        }
562    }
563
564    impl PtxParser for Multicast {
565        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
566            // Try MulticastCluster
567            {
568                let saved_pos = stream.position();
569                if stream.expect_string(".multicast::cluster").is_ok() {
570                    return Ok(Multicast::MulticastCluster);
571                }
572                stream.set_position(saved_pos);
573            }
574            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
575            let expected = &[".multicast::cluster"];
576            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
577            Err(crate::parser::unexpected_value(span, expected, found))
578        }
579    }
580
581    impl PtxParser for Src {
582        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
583            // Try Global
584            {
585                let saved_pos = stream.position();
586                if stream.expect_string(".global").is_ok() {
587                    return Ok(Src::Global);
588                }
589                stream.set_position(saved_pos);
590            }
591            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
592            let expected = &[".global"];
593            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
594            Err(crate::parser::unexpected_value(span, expected, found))
595        }
596    }
597
598    impl PtxParser for CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismMulticastCtaGroupLevelCacheHint {
599        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
600            stream.expect_string("cp")?;
601            stream.expect_string(".async")?;
602            let async_ = ();
603            stream.expect_complete()?;
604            stream.expect_string(".bulk")?;
605            let bulk = ();
606            stream.expect_complete()?;
607            stream.expect_string(".tensor")?;
608            let tensor = ();
609            stream.expect_complete()?;
610            let dim = Dim::parse(stream)?;
611            stream.expect_complete()?;
612            let dst = Dst::parse(stream)?;
613            stream.expect_complete()?;
614            let src = Src::parse(stream)?;
615            stream.expect_complete()?;
616            let saved_pos = stream.position();
617            let load_mode = match LoadMode::parse(stream) {
618                Ok(val) => Some(val),
619                Err(_) => {
620                    stream.set_position(saved_pos);
621                    None
622                }
623            };
624            stream.expect_complete()?;
625            let completion_mechanism = CompletionMechanism::parse(stream)?;
626            stream.expect_complete()?;
627            let saved_pos = stream.position();
628            let multicast = match Multicast::parse(stream) {
629                Ok(val) => Some(val),
630                Err(_) => {
631                    stream.set_position(saved_pos);
632                    None
633                }
634            };
635            stream.expect_complete()?;
636            let saved_pos = stream.position();
637            let cta_group = match CtaGroup::parse(stream) {
638                Ok(val) => Some(val),
639                Err(_) => {
640                    stream.set_position(saved_pos);
641                    None
642                }
643            };
644            stream.expect_complete()?;
645            let saved_pos = stream.position();
646            let level_cache_hint = match LevelCacheHint::parse(stream) {
647                Ok(val) => Some(val),
648                Err(_) => {
649                    stream.set_position(saved_pos);
650                    None
651                }
652            };
653            stream.expect_complete()?;
654            let dstmem = AddressOperand::parse(stream)?;
655            stream.expect_complete()?;
656            stream.expect(&PtxToken::Comma)?;
657            let tensormap = TexHandler2::parse(stream)?;
658            stream.expect_complete()?;
659            stream.expect(&PtxToken::Comma)?;
660            let mbar = AddressOperand::parse(stream)?;
661            stream.expect_complete()?;
662            let saved_pos = stream.position();
663            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
664            if !has_comma {
665                stream.set_position(saved_pos);
666            }
667            let saved_pos = stream.position();
668            let im2colinfo = match GeneralOperand::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 has_comma = stream.expect(&PtxToken::Comma).is_ok();
678            if !has_comma {
679                stream.set_position(saved_pos);
680            }
681            let saved_pos = stream.position();
682            let ctamask = match GeneralOperand::parse(stream) {
683                Ok(val) => Some(val),
684                Err(_) => {
685                    stream.set_position(saved_pos);
686                    None
687                }
688            };
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            stream.expect_complete()?;
705            stream.expect(&PtxToken::Semicolon)?;
706            Ok(CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismMulticastCtaGroupLevelCacheHint {
707                async_,
708                bulk,
709                tensor,
710                dim,
711                dst,
712                src,
713                load_mode,
714                completion_mechanism,
715                multicast,
716                cta_group,
717                level_cache_hint,
718                dstmem,
719                tensormap,
720                mbar,
721                im2colinfo,
722                ctamask,
723                cache_policy,
724            })
725        }
726    }
727
728
729}
730
731pub mod section_2 {
732    use super::*;
733    use crate::r#type::instruction::cp_async_bulk_tensor::section_2::*;
734
735    // ============================================================================
736    // Generated enum parsers
737    // ============================================================================
738
739    impl PtxParser for CompletionMechanism {
740        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
741            // Try BulkGroup
742            {
743                let saved_pos = stream.position();
744                if stream.expect_string(".bulk_group").is_ok() {
745                    return Ok(CompletionMechanism::BulkGroup);
746                }
747                stream.set_position(saved_pos);
748            }
749            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
750            let expected = &[".bulk_group"];
751            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
752            Err(crate::parser::unexpected_value(span, expected, found))
753        }
754    }
755
756    impl PtxParser for Dim {
757        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
758            // Try _1d
759            {
760                let saved_pos = stream.position();
761                if stream.expect_string(".1d").is_ok() {
762                    return Ok(Dim::_1d);
763                }
764                stream.set_position(saved_pos);
765            }
766            let saved_pos = stream.position();
767            // Try _2d
768            {
769                let saved_pos = stream.position();
770                if stream.expect_string(".2d").is_ok() {
771                    return Ok(Dim::_2d);
772                }
773                stream.set_position(saved_pos);
774            }
775            stream.set_position(saved_pos);
776            let saved_pos = stream.position();
777            // Try _3d
778            {
779                let saved_pos = stream.position();
780                if stream.expect_string(".3d").is_ok() {
781                    return Ok(Dim::_3d);
782                }
783                stream.set_position(saved_pos);
784            }
785            stream.set_position(saved_pos);
786            let saved_pos = stream.position();
787            // Try _4d
788            {
789                let saved_pos = stream.position();
790                if stream.expect_string(".4d").is_ok() {
791                    return Ok(Dim::_4d);
792                }
793                stream.set_position(saved_pos);
794            }
795            stream.set_position(saved_pos);
796            let saved_pos = stream.position();
797            // Try _5d
798            {
799                let saved_pos = stream.position();
800                if stream.expect_string(".5d").is_ok() {
801                    return Ok(Dim::_5d);
802                }
803                stream.set_position(saved_pos);
804            }
805            stream.set_position(saved_pos);
806            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
807            let expected = &[".1d", ".2d", ".3d", ".4d", ".5d"];
808            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
809            Err(crate::parser::unexpected_value(span, expected, found))
810        }
811    }
812
813    impl PtxParser for Dst {
814        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
815            // Try Global
816            {
817                let saved_pos = stream.position();
818                if stream.expect_string(".global").is_ok() {
819                    return Ok(Dst::Global);
820                }
821                stream.set_position(saved_pos);
822            }
823            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
824            let expected = &[".global"];
825            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
826            Err(crate::parser::unexpected_value(span, expected, found))
827        }
828    }
829
830    impl PtxParser for LevelCacheHint {
831        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
832            // Try L2CacheHint
833            {
834                let saved_pos = stream.position();
835                if stream.expect_string(".L2::cache_hint").is_ok() {
836                    return Ok(LevelCacheHint::L2CacheHint);
837                }
838                stream.set_position(saved_pos);
839            }
840            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
841            let expected = &[".L2::cache_hint"];
842            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
843            Err(crate::parser::unexpected_value(span, expected, found))
844        }
845    }
846
847    impl PtxParser for LoadMode {
848        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
849            // Try TileScatter4
850            {
851                let saved_pos = stream.position();
852                if stream.expect_string(".tile::scatter4").is_ok() {
853                    return Ok(LoadMode::TileScatter4);
854                }
855                stream.set_position(saved_pos);
856            }
857            let saved_pos = stream.position();
858            // Try Im2colNoOffs
859            {
860                let saved_pos = stream.position();
861                if stream.expect_string(".im2col_no_offs").is_ok() {
862                    return Ok(LoadMode::Im2colNoOffs);
863                }
864                stream.set_position(saved_pos);
865            }
866            stream.set_position(saved_pos);
867            let saved_pos = stream.position();
868            // Try Tile
869            {
870                let saved_pos = stream.position();
871                if stream.expect_string(".tile").is_ok() {
872                    return Ok(LoadMode::Tile);
873                }
874                stream.set_position(saved_pos);
875            }
876            stream.set_position(saved_pos);
877            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
878            let expected = &[".tile::scatter4", ".im2col_no_offs", ".tile"];
879            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
880            Err(crate::parser::unexpected_value(span, expected, found))
881        }
882    }
883
884    impl PtxParser for Src {
885        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
886            // Try SharedCta
887            {
888                let saved_pos = stream.position();
889                if stream.expect_string(".shared::cta").is_ok() {
890                    return Ok(Src::SharedCta);
891                }
892                stream.set_position(saved_pos);
893            }
894            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
895            let expected = &[".shared::cta"];
896            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
897            Err(crate::parser::unexpected_value(span, expected, found))
898        }
899    }
900
901    impl PtxParser for CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismLevelCacheHint {
902        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
903            stream.expect_string("cp")?;
904            stream.expect_string(".async")?;
905            let async_ = ();
906            stream.expect_complete()?;
907            stream.expect_string(".bulk")?;
908            let bulk = ();
909            stream.expect_complete()?;
910            stream.expect_string(".tensor")?;
911            let tensor = ();
912            stream.expect_complete()?;
913            let dim = Dim::parse(stream)?;
914            stream.expect_complete()?;
915            let dst = Dst::parse(stream)?;
916            stream.expect_complete()?;
917            let src = Src::parse(stream)?;
918            stream.expect_complete()?;
919            let saved_pos = stream.position();
920            let load_mode = match LoadMode::parse(stream) {
921                Ok(val) => Some(val),
922                Err(_) => {
923                    stream.set_position(saved_pos);
924                    None
925                }
926            };
927            stream.expect_complete()?;
928            let completion_mechanism = CompletionMechanism::parse(stream)?;
929            stream.expect_complete()?;
930            let saved_pos = stream.position();
931            let level_cache_hint = match LevelCacheHint::parse(stream) {
932                Ok(val) => Some(val),
933                Err(_) => {
934                    stream.set_position(saved_pos);
935                    None
936                }
937            };
938            stream.expect_complete()?;
939            let tensormap = TexHandler2::parse(stream)?;
940            stream.expect_complete()?;
941            stream.expect(&PtxToken::Comma)?;
942            let srcmem = AddressOperand::parse(stream)?;
943            stream.expect_complete()?;
944            let saved_pos = stream.position();
945            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
946            if !has_comma {
947                stream.set_position(saved_pos);
948            }
949            let saved_pos = stream.position();
950            let cache_policy = match GeneralOperand::parse(stream) {
951                Ok(val) => Some(val),
952                Err(_) => {
953                    stream.set_position(saved_pos);
954                    None
955                }
956            };
957            stream.expect_complete()?;
958            stream.expect_complete()?;
959            stream.expect(&PtxToken::Semicolon)?;
960            Ok(CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismLevelCacheHint {
961                async_,
962                bulk,
963                tensor,
964                dim,
965                dst,
966                src,
967                load_mode,
968                completion_mechanism,
969                level_cache_hint,
970                tensormap,
971                srcmem,
972                cache_policy,
973            })
974        }
975    }
976
977
978}
979
ptx_parser/parser/instruction/cp_async_bulk_tensor.rs

ptx_parser/parser/instruction/
cp_async_bulk_tensor.rs
