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            push_opcode(tokens, "tcgen05");
21            push_directive(tokens, "commit");
22            match &self.cta_group {
23                CtaGroup::CtaGroup1 => {
24                    push_directive(tokens, "cta_group::1");
25                }
26                CtaGroup::CtaGroup2 => {
27                    push_directive(tokens, "cta_group::2");
28                }
29            }
30            match &self.completion_mechanism {
31                CompletionMechanism::MbarrierArriveOne => {
32                    push_directive(tokens, "mbarrier::arrive::one");
33                }
34            }
35            if self.shared_cluster {
36                push_directive(tokens, "shared::cluster");
37            }
38            if let Some(multicast_0) = self.multicast.as_ref() {
39                match multicast_0 {
40                    Multicast::MulticastCluster => {
41                        push_directive(tokens, "multicast::cluster");
42                    }
43                }
44            }
45            push_directive(tokens, "b64");
46            self.mbar.unparse_tokens(tokens);
47            if self.ctamask.is_some() {
48                tokens.push(PtxToken::Comma);
49            }
50            if let Some(opt_1) = self.ctamask.as_ref() {
51                opt_1.unparse_tokens(tokens);
52            }
53            tokens.push(PtxToken::Semicolon);
54        }
55    }
56}
ptx_parser/unparser/instruction/tcgen05_commit.rs

ptx_parser/unparser/instruction/
tcgen05_commit.rs
