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
110pub mod section_1 {
111    use super::*;
112    use crate::r#type::instruction::st_async::section_1::*;
113
114    impl PtxUnparser for StAsyncMmioSemScopeSsType {
115        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
116            push_opcode(tokens, "st");
117            push_directive(tokens, "async");
118            if self.mmio {
119                push_directive(tokens, "mmio");
120            }
121            match &self.sem {
122                Sem::Release => {
123                    push_directive(tokens, "release");
124                }
125            }
126            match &self.scope {
127                Scope::Gpu => {
128                    push_directive(tokens, "gpu");
129                }
130                Scope::Sys => {
131                    push_directive(tokens, "sys");
132                }
133            }
134            if let Some(ss_5) = self.ss.as_ref() {
135                match ss_5 {
136                    Ss::Global => {
137                        push_directive(tokens, "global");
138                    }
139                }
140            }
141            match &self.type_ {
142                Type::B16 => {
143                    push_directive(tokens, "b16");
144                }
145                Type::B32 => {
146                    push_directive(tokens, "b32");
147                }
148                Type::B64 => {
149                    push_directive(tokens, "b64");
150                }
151                Type::U16 => {
152                    push_directive(tokens, "u16");
153                }
154                Type::U32 => {
155                    push_directive(tokens, "u32");
156                }
157                Type::U64 => {
158                    push_directive(tokens, "u64");
159                }
160                Type::S16 => {
161                    push_directive(tokens, "s16");
162                }
163                Type::S32 => {
164                    push_directive(tokens, "s32");
165                }
166                Type::S64 => {
167                    push_directive(tokens, "s64");
168                }
169                Type::F32 => {
170                    push_directive(tokens, "f32");
171                }
172                Type::F64 => {
173                    push_directive(tokens, "f64");
174                }
175                Type::B8 => {
176                    push_directive(tokens, "b8");
177                }
178                Type::U8 => {
179                    push_directive(tokens, "u8");
180                }
181                Type::S8 => {
182                    push_directive(tokens, "s8");
183                }
184            }
185            self.a.unparse_tokens(tokens);
186            tokens.push(PtxToken::Comma);
187            self.b.unparse_tokens(tokens);
188            tokens.push(PtxToken::Semicolon);
189        }
190    }
191}
ptx_parser/unparser/instruction/st_async.rs

ptx_parser/unparser/instruction/
st_async.rs
