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

ptx_parser/parser/instruction/
cp_reduce_async_bulk.rs
