Skip to main content

ptx_parser/unparser/instruction/
tcgen05_fence.rs

1//! Original PTX specification:
2//!
3//! tcgen05.fence::before_thread_sync ;
4//! tcgen05.fence::after_thread_sync  ;
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::tcgen05_fence::section_0::*;
14
15    impl PtxUnparser for Tcgen05FenceBeforeThreadSync {
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, "tcgen05");
21            push_directive(tokens, "fence::before_thread_sync");
22            tokens.push(PtxToken::Semicolon);
23            if spaced {
24                tokens.push(PtxToken::Newline);
25            }
26        }
27    }
28
29    impl PtxUnparser for Tcgen05FenceAfterThreadSync {
30        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
31            self.unparse_tokens_mode(tokens, false);
32        }
33        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
34            push_opcode(tokens, "tcgen05");
35            push_directive(tokens, "fence::after_thread_sync");
36            tokens.push(PtxToken::Semicolon);
37            if spaced {
38                tokens.push(PtxToken::Newline);
39            }
40        }
41    }
42}