ptx_parser/unparser/instruction/
tcgen05_commit.rs

1//! Original PTX specification:
2//!
3//! tcgen05.commit.cta_group.completion_mechanism{.shared::cluster}{.multicast}.b64
4//! [mbar] {, ctaMask};
5//! .completion_mechanism = { .mbarrier::arrive::one };
6//! .cta_group            = { .cta_group::1, .cta_group::2 };
7//! .multicast            = { .multicast::cluster };
8
9#![allow(unused)]
10
11use crate::lexer::PtxToken;
12use crate::unparser::{PtxUnparser, common::*};
13
14pub mod section_0 {
15    use super::*;
16    use crate::r#type::instruction::tcgen05_commit::section_0::*;
17
18    impl PtxUnparser for Tcgen05CommitCtaGroupCompletionMechanismSharedClusterMulticastB64 {
19        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
20            self.unparse_tokens_mode(tokens, false);
21        }
22        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
23            push_opcode(tokens, "tcgen05");
24            push_directive(tokens, "commit");
25            match &self.cta_group {
26                CtaGroup::CtaGroup1 => {
27                    push_directive(tokens, "cta_group::1");
28                }
29                CtaGroup::CtaGroup2 => {
30                    push_directive(tokens, "cta_group::2");
31                }
32            }
33            match &self.completion_mechanism {
34                CompletionMechanism::MbarrierArriveOne => {
35                    push_directive(tokens, "mbarrier::arrive::one");
36                }
37            }
38            if self.shared_cluster {
39                push_directive(tokens, "shared::cluster");
40            }
41            if let Some(multicast_0) = self.multicast.as_ref() {
42                match multicast_0 {
43                    Multicast::MulticastCluster => {
44                        push_directive(tokens, "multicast::cluster");
45                    }
46                }
47            }
48            push_directive(tokens, "b64");
49            if spaced {
50                tokens.push(PtxToken::Space);
51            }
52            self.mbar.unparse_tokens_mode(tokens, spaced);
53            if self.ctamask.is_some() {
54                tokens.push(PtxToken::Comma);
55            }
56            if let Some(opt_1) = self.ctamask.as_ref() {
57                if spaced {
58                    tokens.push(PtxToken::Space);
59                }
60                opt_1.unparse_tokens_mode(tokens, spaced);
61            }
62            tokens.push(PtxToken::Semicolon);
63            if spaced {
64                tokens.push(PtxToken::Newline);
65            }
66        }
67    }
68}
ptx_parser/unparser/instruction/tcgen05_commit.rs

ptx_parser/unparser/instruction/
tcgen05_commit.rs
