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::lexer::PtxToken;
8use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
9use crate::r#type::common::*;
10
11pub mod section_0 {
12    use super::*;
13    use crate::r#type::instruction::cp_async_bulk_wait_group::section_0::*;
14
15    impl PtxParser for CpAsyncBulkWaitGroupRead {
16        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
17            stream.expect_string("cp")?;
18            stream.expect_string(".async")?;
19            let async_ = ();
20            stream.expect_complete()?;
21            stream.expect_string(".bulk")?;
22            let bulk = ();
23            stream.expect_complete()?;
24            stream.expect_string(".wait_group")?;
25            let wait_group = ();
26            stream.expect_complete()?;
27            let saved_pos = stream.position();
28            let read = stream.expect_string(".read").is_ok();
29            if !read {
30                stream.set_position(saved_pos);
31            }
32            stream.expect_complete()?;
33            let n = GeneralOperand::parse(stream)?;
34            stream.expect_complete()?;
35            stream.expect_complete()?;
36            stream.expect(&PtxToken::Semicolon)?;
37            Ok(CpAsyncBulkWaitGroupRead {
38                async_,
39                bulk,
40                wait_group,
41                read,
42                n,
43            })
44        }
45    }
46}

ptx_parser/parser/instruction/cp_async_bulk_wait_group.rs

ptx_parser/parser/instruction/
cp_async_bulk_wait_group.rs