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