ptx_parser/unparser/instruction/tcgen05_st.rs
1//! Original PTX specification:
2//!
3//! tcgen05.st.sync.aligned.shape1.num{.unpack}.b32 [taddr], r;
4//! tcgen05.st.sync.aligned.shape2.num{.unpack}.b32 [taddr], immHalfSplitoff, r;
5//! .shape1 = { .16x64b, .16x128b, .16x256b, .32x32b };
6//! .shape2 = { .16x32bx2 };
7//! .num = { .x1, .x2, .x4, .x8, .x16, .x32, .x64, .x128 };
8//! .unpack = { .unpack::16b };
9
10#![allow(unused)]
11
12use crate::lexer::PtxToken;
13use crate::unparser::{PtxUnparser, common::*};
14
15pub mod section_0 {
16 use super::*;
17 use crate::r#type::instruction::tcgen05_st::section_0::*;
18
19 impl PtxUnparser for Tcgen05StSyncAlignedShape1NumUnpackB32 {
20 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
21 push_opcode(tokens, "tcgen05");
22 push_directive(tokens, "st");
23 push_directive(tokens, "sync");
24 push_directive(tokens, "aligned");
25 match &self.shape1 {
26 Shape1::_16x128b => {
27 push_directive(tokens, "16x128b");
28 }
29 Shape1::_16x256b => {
30 push_directive(tokens, "16x256b");
31 }
32 Shape1::_16x64b => {
33 push_directive(tokens, "16x64b");
34 }
35 Shape1::_32x32b => {
36 push_directive(tokens, "32x32b");
37 }
38 }
39 match &self.num {
40 Num::X128 => {
41 push_directive(tokens, "x128");
42 }
43 Num::X16 => {
44 push_directive(tokens, "x16");
45 }
46 Num::X32 => {
47 push_directive(tokens, "x32");
48 }
49 Num::X64 => {
50 push_directive(tokens, "x64");
51 }
52 Num::X1 => {
53 push_directive(tokens, "x1");
54 }
55 Num::X2 => {
56 push_directive(tokens, "x2");
57 }
58 Num::X4 => {
59 push_directive(tokens, "x4");
60 }
61 Num::X8 => {
62 push_directive(tokens, "x8");
63 }
64 }
65 if let Some(unpack_0) = self.unpack.as_ref() {
66 match unpack_0 {
67 Unpack::Unpack16b => {
68 push_directive(tokens, "unpack::16b");
69 }
70 }
71 }
72 push_directive(tokens, "b32");
73 self.taddr.unparse_tokens(tokens);
74 tokens.push(PtxToken::Comma);
75 self.r.unparse_tokens(tokens);
76 tokens.push(PtxToken::Semicolon);
77 }
78 }
79
80 impl PtxUnparser for Tcgen05StSyncAlignedShape2NumUnpackB32 {
81 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
82 push_opcode(tokens, "tcgen05");
83 push_directive(tokens, "st");
84 push_directive(tokens, "sync");
85 push_directive(tokens, "aligned");
86 match &self.shape2 {
87 Shape2::_16x32bx2 => {
88 push_directive(tokens, "16x32bx2");
89 }
90 }
91 match &self.num {
92 Num::X128 => {
93 push_directive(tokens, "x128");
94 }
95 Num::X16 => {
96 push_directive(tokens, "x16");
97 }
98 Num::X32 => {
99 push_directive(tokens, "x32");
100 }
101 Num::X64 => {
102 push_directive(tokens, "x64");
103 }
104 Num::X1 => {
105 push_directive(tokens, "x1");
106 }
107 Num::X2 => {
108 push_directive(tokens, "x2");
109 }
110 Num::X4 => {
111 push_directive(tokens, "x4");
112 }
113 Num::X8 => {
114 push_directive(tokens, "x8");
115 }
116 }
117 if let Some(unpack_1) = self.unpack.as_ref() {
118 match unpack_1 {
119 Unpack::Unpack16b => {
120 push_directive(tokens, "unpack::16b");
121 }
122 }
123 }
124 push_directive(tokens, "b32");
125 self.taddr.unparse_tokens(tokens);
126 tokens.push(PtxToken::Comma);
127 self.immhalfsplitoff.unparse_tokens(tokens);
128 tokens.push(PtxToken::Comma);
129 self.r.unparse_tokens(tokens);
130 tokens.push(PtxToken::Semicolon);
131 }
132 }
133
134}
135