ptx_parser/unparser/instruction/
mbarrier_test_wait.rs

1//! Original PTX specification:
2//!
3//! mbarrier.test_wait{.sem}{.scope}{.state}.b64        waitComplete, [addr], state;
4//! mbarrier.test_wait.parity{.sem}{.scope}{.state}.b64 waitComplete, [addr], phaseParity;
5//! mbarrier.try_wait{.sem}{.scope}{.state}.b64         waitComplete, [addr], state {, suspendTimeHint};
6//! mbarrier.try_wait.parity{.sem}{.scope}{.state}.b64  waitComplete, [addr], phaseParity {, suspendTimeHint};
7//! .sem   = { .acquire, .relaxed };
8//! .scope = { .cta, .cluster };
9//! .state = { .shared, .shared::cta}
10
11#![allow(unused)]
12
13use crate::lexer::PtxToken;
14use crate::unparser::{PtxUnparser, common::*};
15
16pub mod section_0 {
17    use super::*;
18    use crate::r#type::instruction::mbarrier_test_wait::section_0::*;
19
20    impl PtxUnparser for MbarrierTestWaitSemScopeStateB64 {
21        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
22            self.unparse_tokens_mode(tokens, false);
23        }
24        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
25            push_opcode(tokens, "mbarrier");
26            push_directive(tokens, "test_wait");
27            if let Some(sem_0) = self.sem.as_ref() {
28                match sem_0 {
29                    Sem::Acquire => {
30                        push_directive(tokens, "acquire");
31                    }
32                    Sem::Relaxed => {
33                        push_directive(tokens, "relaxed");
34                    }
35                }
36            }
37            if let Some(scope_1) = self.scope.as_ref() {
38                match scope_1 {
39                    Scope::Cluster => {
40                        push_directive(tokens, "cluster");
41                    }
42                    Scope::Cta => {
43                        push_directive(tokens, "cta");
44                    }
45                }
46            }
47            if let Some(state_2) = self.state.as_ref() {
48                match state_2 {
49                    State::SharedCta => {
50                        push_directive(tokens, "shared::cta");
51                    }
52                    State::Shared => {
53                        push_directive(tokens, "shared");
54                    }
55                }
56            }
57            push_directive(tokens, "b64");
58            if spaced {
59                tokens.push(PtxToken::Space);
60            }
61            self.waitcomplete.unparse_tokens_mode(tokens, spaced);
62            tokens.push(PtxToken::Comma);
63            if spaced {
64                tokens.push(PtxToken::Space);
65            }
66            self.addr.unparse_tokens_mode(tokens, spaced);
67            tokens.push(PtxToken::Comma);
68            if spaced {
69                tokens.push(PtxToken::Space);
70            }
71            self.state2.unparse_tokens_mode(tokens, spaced);
72            tokens.push(PtxToken::Semicolon);
73            if spaced {
74                tokens.push(PtxToken::Newline);
75            }
76        }
77    }
78
79    impl PtxUnparser for MbarrierTestWaitParitySemScopeStateB64 {
80        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
81            self.unparse_tokens_mode(tokens, false);
82        }
83        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
84            push_opcode(tokens, "mbarrier");
85            push_directive(tokens, "test_wait");
86            push_directive(tokens, "parity");
87            if let Some(sem_3) = self.sem.as_ref() {
88                match sem_3 {
89                    Sem::Acquire => {
90                        push_directive(tokens, "acquire");
91                    }
92                    Sem::Relaxed => {
93                        push_directive(tokens, "relaxed");
94                    }
95                }
96            }
97            if let Some(scope_4) = self.scope.as_ref() {
98                match scope_4 {
99                    Scope::Cluster => {
100                        push_directive(tokens, "cluster");
101                    }
102                    Scope::Cta => {
103                        push_directive(tokens, "cta");
104                    }
105                }
106            }
107            if let Some(state_5) = self.state.as_ref() {
108                match state_5 {
109                    State::SharedCta => {
110                        push_directive(tokens, "shared::cta");
111                    }
112                    State::Shared => {
113                        push_directive(tokens, "shared");
114                    }
115                }
116            }
117            push_directive(tokens, "b64");
118            if spaced {
119                tokens.push(PtxToken::Space);
120            }
121            self.waitcomplete.unparse_tokens_mode(tokens, spaced);
122            tokens.push(PtxToken::Comma);
123            if spaced {
124                tokens.push(PtxToken::Space);
125            }
126            self.addr.unparse_tokens_mode(tokens, spaced);
127            tokens.push(PtxToken::Comma);
128            if spaced {
129                tokens.push(PtxToken::Space);
130            }
131            self.phaseparity.unparse_tokens_mode(tokens, spaced);
132            tokens.push(PtxToken::Semicolon);
133            if spaced {
134                tokens.push(PtxToken::Newline);
135            }
136        }
137    }
138
139    impl PtxUnparser for MbarrierTryWaitSemScopeStateB64 {
140        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
141            self.unparse_tokens_mode(tokens, false);
142        }
143        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
144            push_opcode(tokens, "mbarrier");
145            push_directive(tokens, "try_wait");
146            if let Some(sem_6) = self.sem.as_ref() {
147                match sem_6 {
148                    Sem::Acquire => {
149                        push_directive(tokens, "acquire");
150                    }
151                    Sem::Relaxed => {
152                        push_directive(tokens, "relaxed");
153                    }
154                }
155            }
156            if let Some(scope_7) = self.scope.as_ref() {
157                match scope_7 {
158                    Scope::Cluster => {
159                        push_directive(tokens, "cluster");
160                    }
161                    Scope::Cta => {
162                        push_directive(tokens, "cta");
163                    }
164                }
165            }
166            if let Some(state_8) = self.state.as_ref() {
167                match state_8 {
168                    State::SharedCta => {
169                        push_directive(tokens, "shared::cta");
170                    }
171                    State::Shared => {
172                        push_directive(tokens, "shared");
173                    }
174                }
175            }
176            push_directive(tokens, "b64");
177            if spaced {
178                tokens.push(PtxToken::Space);
179            }
180            self.waitcomplete.unparse_tokens_mode(tokens, spaced);
181            tokens.push(PtxToken::Comma);
182            if spaced {
183                tokens.push(PtxToken::Space);
184            }
185            self.addr.unparse_tokens_mode(tokens, spaced);
186            tokens.push(PtxToken::Comma);
187            if spaced {
188                tokens.push(PtxToken::Space);
189            }
190            self.state2.unparse_tokens_mode(tokens, spaced);
191            if self.suspendtimehint.is_some() {
192                tokens.push(PtxToken::Comma);
193            }
194            if let Some(opt_9) = self.suspendtimehint.as_ref() {
195                if spaced {
196                    tokens.push(PtxToken::Space);
197                }
198                opt_9.unparse_tokens_mode(tokens, spaced);
199            }
200            tokens.push(PtxToken::Semicolon);
201            if spaced {
202                tokens.push(PtxToken::Newline);
203            }
204        }
205    }
206
207    impl PtxUnparser for MbarrierTryWaitParitySemScopeStateB64 {
208        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
209            self.unparse_tokens_mode(tokens, false);
210        }
211        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
212            push_opcode(tokens, "mbarrier");
213            push_directive(tokens, "try_wait");
214            push_directive(tokens, "parity");
215            if let Some(sem_10) = self.sem.as_ref() {
216                match sem_10 {
217                    Sem::Acquire => {
218                        push_directive(tokens, "acquire");
219                    }
220                    Sem::Relaxed => {
221                        push_directive(tokens, "relaxed");
222                    }
223                }
224            }
225            if let Some(scope_11) = self.scope.as_ref() {
226                match scope_11 {
227                    Scope::Cluster => {
228                        push_directive(tokens, "cluster");
229                    }
230                    Scope::Cta => {
231                        push_directive(tokens, "cta");
232                    }
233                }
234            }
235            if let Some(state_12) = self.state.as_ref() {
236                match state_12 {
237                    State::SharedCta => {
238                        push_directive(tokens, "shared::cta");
239                    }
240                    State::Shared => {
241                        push_directive(tokens, "shared");
242                    }
243                }
244            }
245            push_directive(tokens, "b64");
246            if spaced {
247                tokens.push(PtxToken::Space);
248            }
249            self.waitcomplete.unparse_tokens_mode(tokens, spaced);
250            tokens.push(PtxToken::Comma);
251            if spaced {
252                tokens.push(PtxToken::Space);
253            }
254            self.addr.unparse_tokens_mode(tokens, spaced);
255            tokens.push(PtxToken::Comma);
256            if spaced {
257                tokens.push(PtxToken::Space);
258            }
259            self.phaseparity.unparse_tokens_mode(tokens, spaced);
260            if self.suspendtimehint.is_some() {
261                tokens.push(PtxToken::Comma);
262            }
263            if let Some(opt_13) = self.suspendtimehint.as_ref() {
264                if spaced {
265                    tokens.push(PtxToken::Space);
266                }
267                opt_13.unparse_tokens_mode(tokens, spaced);
268            }
269            tokens.push(PtxToken::Semicolon);
270            if spaced {
271                tokens.push(PtxToken::Newline);
272            }
273        }
274    }
275}
ptx_parser/unparser/instruction/mbarrier_test_wait.rs

ptx_parser/unparser/instruction/
mbarrier_test_wait.rs
