Skip to main content

ptx_parser/unparser/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::unparser::{PtxUnparser, common::*};
10
11pub mod section_0 {
12    use super::*;
13    use crate::r#type::instruction::cp_async_wait_group::section_0::*;
14
15    impl PtxUnparser for CpAsyncWaitGroup {
16        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
17            self.unparse_tokens_mode(tokens, false);
18        }
19        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
20            push_opcode(tokens, "cp");
21            push_directive(tokens, "async");
22            push_directive(tokens, "wait_group");
23            if spaced {
24                tokens.push(PtxToken::Space);
25            }
26            self.n.unparse_tokens_mode(tokens, spaced);
27            tokens.push(PtxToken::Semicolon);
28            if spaced {
29                tokens.push(PtxToken::Newline);
30            }
31        }
32    }
33
34    impl PtxUnparser for CpAsyncWaitAll {
35        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
36            self.unparse_tokens_mode(tokens, false);
37        }
38        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
39            push_opcode(tokens, "cp");
40            push_directive(tokens, "async");
41            push_directive(tokens, "wait_all");
42            tokens.push(PtxToken::Semicolon);
43            if spaced {
44                tokens.push(PtxToken::Newline);
45            }
46        }
47    }
48}