ptx_parser/unparser/instruction/
mbarrier_arrive.rs

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

ptx_parser/unparser/instruction/
mbarrier_arrive.rs
