ptx_parser/unparser/instruction/st_async.rs
1//! Original PTX specification:
2//!
3//! st.async{.sem}{.scope}{.ss}{.completion_mechanism}{.vec}.type [a], b, [mbar];
4//! .sem = { .weak };
5//! .scope = { .cluster };
6//! .ss = { .shared::cluster };
7//! .type = { .b32, .b64,
8//! .u32, .u64,
9//! .s32, .s64,
10//! .f32, .f64 };
11//! .vec = { .v2, .v4 };
12//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
13//! ---------------------------------------------------------
14//! st.async{.mmio}.sem.scope{.ss}.type [a], b;
15//! .sem = { .release };
16//! .scope = { .gpu, .sys };
17//! .ss = { .global };
18//! .type = { .b8, .b16, .b32, .b64,
19//! .u8, .u16, .u32, .u64,
20//! .s8, .s16, .s32, .s64,
21//! .f32, .f64 };
22
23#![allow(unused)]
24
25use crate::lexer::PtxToken;
26use crate::unparser::{PtxUnparser, common::*};
27
28pub mod section_0 {
29 use super::*;
30 use crate::r#type::instruction::st_async::section_0::*;
31
32 impl PtxUnparser for StAsyncSemScopeSsCompletionMechanismVecType {
33 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
34 push_opcode(tokens, "st");
35 push_directive(tokens, "async");
36 if let Some(sem_0) = self.sem.as_ref() {
37 match sem_0 {
38 Sem::Weak => {
39 push_directive(tokens, "weak");
40 }
41 }
42 }
43 if let Some(scope_1) = self.scope.as_ref() {
44 match scope_1 {
45 Scope::Cluster => {
46 push_directive(tokens, "cluster");
47 }
48 }
49 }
50 if let Some(ss_2) = self.ss.as_ref() {
51 match ss_2 {
52 Ss::SharedCluster => {
53 push_directive(tokens, "shared::cluster");
54 }
55 }
56 }
57 if let Some(completion_mechanism_3) = self.completion_mechanism.as_ref() {
58 match completion_mechanism_3 {
59 CompletionMechanism::MbarrierCompleteTxBytes => {
60 push_directive(tokens, "mbarrier::complete_tx::bytes");
61 }
62 }
63 }
64 if let Some(vec_4) = self.vec.as_ref() {
65 match vec_4 {
66 Vec::V2 => {
67 push_directive(tokens, "v2");
68 }
69 Vec::V4 => {
70 push_directive(tokens, "v4");
71 }
72 }
73 }
74 match &self.type_ {
75 Type::B32 => {
76 push_directive(tokens, "b32");
77 }
78 Type::B64 => {
79 push_directive(tokens, "b64");
80 }
81 Type::U32 => {
82 push_directive(tokens, "u32");
83 }
84 Type::U64 => {
85 push_directive(tokens, "u64");
86 }
87 Type::S32 => {
88 push_directive(tokens, "s32");
89 }
90 Type::S64 => {
91 push_directive(tokens, "s64");
92 }
93 Type::F32 => {
94 push_directive(tokens, "f32");
95 }
96 Type::F64 => {
97 push_directive(tokens, "f64");
98 }
99 }
100 self.a.unparse_tokens(tokens);
101 tokens.push(PtxToken::Comma);
102 self.b.unparse_tokens(tokens);
103 tokens.push(PtxToken::Comma);
104 self.mbar.unparse_tokens(tokens);
105 tokens.push(PtxToken::Semicolon);
106 }
107 }
108
109}
110
111pub mod section_1 {
112 use super::*;
113 use crate::r#type::instruction::st_async::section_1::*;
114
115 impl PtxUnparser for StAsyncMmioSemScopeSsType {
116 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
117 push_opcode(tokens, "st");
118 push_directive(tokens, "async");
119 if self.mmio {
120 push_directive(tokens, "mmio");
121 }
122 match &self.sem {
123 Sem::Release => {
124 push_directive(tokens, "release");
125 }
126 }
127 match &self.scope {
128 Scope::Gpu => {
129 push_directive(tokens, "gpu");
130 }
131 Scope::Sys => {
132 push_directive(tokens, "sys");
133 }
134 }
135 if let Some(ss_5) = self.ss.as_ref() {
136 match ss_5 {
137 Ss::Global => {
138 push_directive(tokens, "global");
139 }
140 }
141 }
142 match &self.type_ {
143 Type::B16 => {
144 push_directive(tokens, "b16");
145 }
146 Type::B32 => {
147 push_directive(tokens, "b32");
148 }
149 Type::B64 => {
150 push_directive(tokens, "b64");
151 }
152 Type::U16 => {
153 push_directive(tokens, "u16");
154 }
155 Type::U32 => {
156 push_directive(tokens, "u32");
157 }
158 Type::U64 => {
159 push_directive(tokens, "u64");
160 }
161 Type::S16 => {
162 push_directive(tokens, "s16");
163 }
164 Type::S32 => {
165 push_directive(tokens, "s32");
166 }
167 Type::S64 => {
168 push_directive(tokens, "s64");
169 }
170 Type::F32 => {
171 push_directive(tokens, "f32");
172 }
173 Type::F64 => {
174 push_directive(tokens, "f64");
175 }
176 Type::B8 => {
177 push_directive(tokens, "b8");
178 }
179 Type::U8 => {
180 push_directive(tokens, "u8");
181 }
182 Type::S8 => {
183 push_directive(tokens, "s8");
184 }
185 }
186 self.a.unparse_tokens(tokens);
187 tokens.push(PtxToken::Comma);
188 self.b.unparse_tokens(tokens);
189 tokens.push(PtxToken::Semicolon);
190 }
191 }
192
193}
194