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            self.unparse_tokens_mode(tokens, false);
35        }
36        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
37            push_opcode(tokens, "st");
38            push_directive(tokens, "async");
39            if let Some(sem_0) = self.sem.as_ref() {
40                match sem_0 {
41                    Sem::Weak => {
42                        push_directive(tokens, "weak");
43                    }
44                }
45            }
46            if let Some(scope_1) = self.scope.as_ref() {
47                match scope_1 {
48                    Scope::Cluster => {
49                        push_directive(tokens, "cluster");
50                    }
51                }
52            }
53            if let Some(ss_2) = self.ss.as_ref() {
54                match ss_2 {
55                    Ss::SharedCluster => {
56                        push_directive(tokens, "shared::cluster");
57                    }
58                }
59            }
60            if let Some(completion_mechanism_3) = self.completion_mechanism.as_ref() {
61                match completion_mechanism_3 {
62                    CompletionMechanism::MbarrierCompleteTxBytes => {
63                        push_directive(tokens, "mbarrier::complete_tx::bytes");
64                    }
65                }
66            }
67            if let Some(vec_4) = self.vec.as_ref() {
68                match vec_4 {
69                    Vec::V2 => {
70                        push_directive(tokens, "v2");
71                    }
72                    Vec::V4 => {
73                        push_directive(tokens, "v4");
74                    }
75                }
76            }
77            match &self.type_ {
78                Type::B32 => {
79                    push_directive(tokens, "b32");
80                }
81                Type::B64 => {
82                    push_directive(tokens, "b64");
83                }
84                Type::U32 => {
85                    push_directive(tokens, "u32");
86                }
87                Type::U64 => {
88                    push_directive(tokens, "u64");
89                }
90                Type::S32 => {
91                    push_directive(tokens, "s32");
92                }
93                Type::S64 => {
94                    push_directive(tokens, "s64");
95                }
96                Type::F32 => {
97                    push_directive(tokens, "f32");
98                }
99                Type::F64 => {
100                    push_directive(tokens, "f64");
101                }
102            }
103            if spaced {
104                tokens.push(PtxToken::Space);
105            }
106            self.a.unparse_tokens_mode(tokens, spaced);
107            tokens.push(PtxToken::Comma);
108            if spaced {
109                tokens.push(PtxToken::Space);
110            }
111            self.b.unparse_tokens_mode(tokens, spaced);
112            tokens.push(PtxToken::Comma);
113            if spaced {
114                tokens.push(PtxToken::Space);
115            }
116            self.mbar.unparse_tokens_mode(tokens, spaced);
117            tokens.push(PtxToken::Semicolon);
118            if spaced {
119                tokens.push(PtxToken::Newline);
120            }
121        }
122    }
123}
124
125pub mod section_1 {
126    use super::*;
127    use crate::r#type::instruction::st_async::section_1::*;
128
129    impl PtxUnparser for StAsyncMmioSemScopeSsType {
130        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
131            self.unparse_tokens_mode(tokens, false);
132        }
133        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
134            push_opcode(tokens, "st");
135            push_directive(tokens, "async");
136            if self.mmio {
137                push_directive(tokens, "mmio");
138            }
139            match &self.sem {
140                Sem::Release => {
141                    push_directive(tokens, "release");
142                }
143            }
144            match &self.scope {
145                Scope::Gpu => {
146                    push_directive(tokens, "gpu");
147                }
148                Scope::Sys => {
149                    push_directive(tokens, "sys");
150                }
151            }
152            if let Some(ss_5) = self.ss.as_ref() {
153                match ss_5 {
154                    Ss::Global => {
155                        push_directive(tokens, "global");
156                    }
157                }
158            }
159            match &self.type_ {
160                Type::B16 => {
161                    push_directive(tokens, "b16");
162                }
163                Type::B32 => {
164                    push_directive(tokens, "b32");
165                }
166                Type::B64 => {
167                    push_directive(tokens, "b64");
168                }
169                Type::U16 => {
170                    push_directive(tokens, "u16");
171                }
172                Type::U32 => {
173                    push_directive(tokens, "u32");
174                }
175                Type::U64 => {
176                    push_directive(tokens, "u64");
177                }
178                Type::S16 => {
179                    push_directive(tokens, "s16");
180                }
181                Type::S32 => {
182                    push_directive(tokens, "s32");
183                }
184                Type::S64 => {
185                    push_directive(tokens, "s64");
186                }
187                Type::F32 => {
188                    push_directive(tokens, "f32");
189                }
190                Type::F64 => {
191                    push_directive(tokens, "f64");
192                }
193                Type::B8 => {
194                    push_directive(tokens, "b8");
195                }
196                Type::U8 => {
197                    push_directive(tokens, "u8");
198                }
199                Type::S8 => {
200                    push_directive(tokens, "s8");
201                }
202            }
203            if spaced {
204                tokens.push(PtxToken::Space);
205            }
206            self.a.unparse_tokens_mode(tokens, spaced);
207            tokens.push(PtxToken::Comma);
208            if spaced {
209                tokens.push(PtxToken::Space);
210            }
211            self.b.unparse_tokens_mode(tokens, spaced);
212            tokens.push(PtxToken::Semicolon);
213            if spaced {
214                tokens.push(PtxToken::Newline);
215            }
216        }
217    }
218}
ptx_parser/unparser/instruction/st_async.rs

ptx_parser/unparser/instruction/
st_async.rs
