Skip to main content

ptx_parser/parser/instruction/
wgmma_wait_group.rs

1//! Original PTX specification:
2//!
3//! wgmma.wait_group.sync.aligned 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::wgmma_wait_group::section_0::*;
20
21    impl PtxParser for WgmmaWaitGroupSyncAligned {
22        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
23            try_map(
24                seq_n!(
25                    string_p("wgmma"),
26                    string_p(".wait_group"),
27                    string_p(".sync"),
28                    string_p(".aligned"),
29                    GeneralOperand::parse(),
30                    semicolon_p()
31                ),
32                |(_, wait_group, sync, aligned, n, _), span| {
33                    ok!(WgmmaWaitGroupSyncAligned {
34                        wait_group = wait_group,
35                        sync = sync,
36                        aligned = aligned,
37                        n = n,
38
39                    })
40                },
41            )
42        }
43    }
44}