ptx_parser/parser/instruction/
cp_async_wait_group.rs

1//! Original PTX specification:
2//!
3//! cp.async.wait_group N;
4//! cp.async.wait_all ;
5
6#![allow(unused)]
7
8use crate::parser::{
9    PtxParseError, PtxParser, PtxTokenStream, Span,
10    util::{
11        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
12        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
13    },
14};
15use crate::r#type::common::*;
16use crate::{alt, ok, seq_n};
17
18pub mod section_0 {
19    use super::*;
20    use crate::r#type::instruction::cp_async_wait_group::section_0::*;
21
22    impl PtxParser for CpAsyncWaitGroup {
23        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
24            try_map(
25                seq_n!(
26                    string_p("cp"),
27                    string_p(".async"),
28                    string_p(".wait_group"),
29                    GeneralOperand::parse(),
30                    semicolon_p()
31                ),
32                |(_, async_, wait_group, n, _), span| {
33                    ok!(CpAsyncWaitGroup {
34                        async_ = async_,
35                        wait_group = wait_group,
36                        n = n,
37
38                    })
39                },
40            )
41        }
42    }
43
44    impl PtxParser for CpAsyncWaitAll {
45        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
46            try_map(
47                seq_n!(
48                    string_p("cp"),
49                    string_p(".async"),
50                    string_p(".wait_all"),
51                    semicolon_p()
52                ),
53                |(_, async_, wait_all, _), span| {
54                    ok!(CpAsyncWaitAll {
55                        async_ = async_,
56                        wait_all = wait_all,
57
58                    })
59                },
60            )
61        }
62    }
63}