ptx_parser/parser/instruction/
cp_async_bulk_wait_group.rs

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

ptx_parser/parser/instruction/cp_async_bulk_wait_group.rs

ptx_parser/parser/instruction/
cp_async_bulk_wait_group.rs