ptx_parser/unparser/instruction/
tcgen05_alloc.rs

1//! Original PTX specification:
2//!
3//! tcgen05.alloc.cta_group.sync.aligned{.shared::cta}.b32  [dst], nCols;
4//! tcgen05.dealloc.cta_group.sync.aligned.b32              taddr, nCols;
5//! tcgen05.relinquish_alloc_permit.cta_group.sync.aligned;
6//! .cta_group = { .cta_group::1, .cta_group::2 };
7
8#![allow(unused)]
9
10use crate::lexer::PtxToken;
11use crate::unparser::{PtxUnparser, common::*};
12
13pub mod section_0 {
14    use super::*;
15    use crate::r#type::instruction::tcgen05_alloc::section_0::*;
16
17    impl PtxUnparser for Tcgen05AllocCtaGroupSyncAlignedSharedCtaB32 {
18        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
19            push_opcode(tokens, "tcgen05");
20            push_directive(tokens, "alloc");
21            match &self.cta_group {
22                CtaGroup::CtaGroup1 => {
23                    push_directive(tokens, "cta_group::1");
24                }
25                CtaGroup::CtaGroup2 => {
26                    push_directive(tokens, "cta_group::2");
27                }
28            }
29            push_directive(tokens, "sync");
30            push_directive(tokens, "aligned");
31            if self.shared_cta {
32                push_directive(tokens, "shared::cta");
33            }
34            push_directive(tokens, "b32");
35            self.dst.unparse_tokens(tokens);
36            tokens.push(PtxToken::Comma);
37            self.ncols.unparse_tokens(tokens);
38            tokens.push(PtxToken::Semicolon);
39        }
40    }
41
42    impl PtxUnparser for Tcgen05DeallocCtaGroupSyncAlignedB32 {
43        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
44            push_opcode(tokens, "tcgen05");
45            push_directive(tokens, "dealloc");
46            match &self.cta_group {
47                CtaGroup::CtaGroup1 => {
48                    push_directive(tokens, "cta_group::1");
49                }
50                CtaGroup::CtaGroup2 => {
51                    push_directive(tokens, "cta_group::2");
52                }
53            }
54            push_directive(tokens, "sync");
55            push_directive(tokens, "aligned");
56            push_directive(tokens, "b32");
57            self.taddr.unparse_tokens(tokens);
58            tokens.push(PtxToken::Comma);
59            self.ncols.unparse_tokens(tokens);
60            tokens.push(PtxToken::Semicolon);
61        }
62    }
63
64    impl PtxUnparser for Tcgen05RelinquishAllocPermitCtaGroupSyncAligned {
65        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
66            push_opcode(tokens, "tcgen05");
67            push_directive(tokens, "relinquish_alloc_permit");
68            match &self.cta_group {
69                CtaGroup::CtaGroup1 => {
70                    push_directive(tokens, "cta_group::1");
71                }
72                CtaGroup::CtaGroup2 => {
73                    push_directive(tokens, "cta_group::2");
74                }
75            }
76            push_directive(tokens, "sync");
77            push_directive(tokens, "aligned");
78            tokens.push(PtxToken::Semicolon);
79        }
80    }
81}
ptx_parser/unparser/instruction/tcgen05_alloc.rs

ptx_parser/unparser/instruction/
tcgen05_alloc.rs
