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() { tokens.push(PtxToken::Comma); }
48                    if let Some(opt_1) = self.ctamask.as_ref() {
49                        opt_1.unparse_tokens(tokens);
50                    }
51            tokens.push(PtxToken::Semicolon);
52        }
53    }
54
55}
56
ptx_parser/unparser/instruction/tcgen05_commit.rs

ptx_parser/unparser/instruction/
tcgen05_commit.rs
