ptx_parser/parser/instruction/
cp_reduce_async_bulk.rs

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

ptx_parser/parser/instruction/
cp_reduce_async_bulk.rs
