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

ptx_parser/parser/instruction/
cp_async_bulk_tensor.rs
