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            self.unparse_tokens_mode(tokens, false);
20        }
21        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
22            push_opcode(tokens, "tcgen05");
23            push_directive(tokens, "alloc");
24            match &self.cta_group {
25                CtaGroup::CtaGroup1 => {
26                    push_directive(tokens, "cta_group::1");
27                }
28                CtaGroup::CtaGroup2 => {
29                    push_directive(tokens, "cta_group::2");
30                }
31            }
32            push_directive(tokens, "sync");
33            push_directive(tokens, "aligned");
34            if self.shared_cta {
35                push_directive(tokens, "shared::cta");
36            }
37            push_directive(tokens, "b32");
38            if spaced {
39                tokens.push(PtxToken::Space);
40            }
41            self.dst.unparse_tokens_mode(tokens, spaced);
42            tokens.push(PtxToken::Comma);
43            if spaced {
44                tokens.push(PtxToken::Space);
45            }
46            self.ncols.unparse_tokens_mode(tokens, spaced);
47            tokens.push(PtxToken::Semicolon);
48            if spaced {
49                tokens.push(PtxToken::Newline);
50            }
51        }
52    }
53
54    impl PtxUnparser for Tcgen05DeallocCtaGroupSyncAlignedB32 {
55        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
56            self.unparse_tokens_mode(tokens, false);
57        }
58        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
59            push_opcode(tokens, "tcgen05");
60            push_directive(tokens, "dealloc");
61            match &self.cta_group {
62                CtaGroup::CtaGroup1 => {
63                    push_directive(tokens, "cta_group::1");
64                }
65                CtaGroup::CtaGroup2 => {
66                    push_directive(tokens, "cta_group::2");
67                }
68            }
69            push_directive(tokens, "sync");
70            push_directive(tokens, "aligned");
71            push_directive(tokens, "b32");
72            if spaced {
73                tokens.push(PtxToken::Space);
74            }
75            self.taddr.unparse_tokens_mode(tokens, spaced);
76            tokens.push(PtxToken::Comma);
77            if spaced {
78                tokens.push(PtxToken::Space);
79            }
80            self.ncols.unparse_tokens_mode(tokens, spaced);
81            tokens.push(PtxToken::Semicolon);
82            if spaced {
83                tokens.push(PtxToken::Newline);
84            }
85        }
86    }
87
88    impl PtxUnparser for Tcgen05RelinquishAllocPermitCtaGroupSyncAligned {
89        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
90            self.unparse_tokens_mode(tokens, false);
91        }
92        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
93            push_opcode(tokens, "tcgen05");
94            push_directive(tokens, "relinquish_alloc_permit");
95            match &self.cta_group {
96                CtaGroup::CtaGroup1 => {
97                    push_directive(tokens, "cta_group::1");
98                }
99                CtaGroup::CtaGroup2 => {
100                    push_directive(tokens, "cta_group::2");
101                }
102            }
103            push_directive(tokens, "sync");
104            push_directive(tokens, "aligned");
105            tokens.push(PtxToken::Semicolon);
106            if spaced {
107                tokens.push(PtxToken::Newline);
108            }
109        }
110    }
111}
ptx_parser/unparser/instruction/tcgen05_alloc.rs

ptx_parser/unparser/instruction/
tcgen05_alloc.rs
