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
82}
83
ptx_parser/unparser/instruction/tcgen05_alloc.rs

ptx_parser/unparser/instruction/
tcgen05_alloc.rs
