ptx_parser/unparser/instruction/
mbarrier_arrive_drop.rs

1//! Original PTX specification:
2//!
3//! mbarrier.arrive_drop{.sem}{.scope}{.state}.b64 state,           [addr]{, count};
4//! mbarrier.arrive_drop{.sem}{.scope}{.shared::cluster}.b64           _,   [addr] {,count};
5//! mbarrier.arrive_drop.expect_tx{.state}{.sem}{.scope}.b64 state, [addr], tx_count;
6//! mbarrier.arrive_drop.expect_tx{.shared::cluster}{.sem}{.scope}.b64   _, [addr], tx_count;
7//! mbarrier.arrive_drop.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_drop::section_0::*;
20
21    impl PtxUnparser for MbarrierArriveDropSemScopeStateB64 {
22        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
23            push_opcode(tokens, "mbarrier");
24            push_directive(tokens, "arrive_drop");
25            if let Some(sem_0) = self.sem.as_ref() {
26                match sem_0 {
27                    Sem::Release => {
28                        push_directive(tokens, "release");
29                    }
30                    Sem::Relaxed => {
31                        push_directive(tokens, "relaxed");
32                    }
33                }
34            }
35            if let Some(scope_1) = self.scope.as_ref() {
36                match scope_1 {
37                    Scope::Cluster => {
38                        push_directive(tokens, "cluster");
39                    }
40                    Scope::Cta => {
41                        push_directive(tokens, "cta");
42                    }
43                }
44            }
45            if let Some(state_2) = self.state.as_ref() {
46                match state_2 {
47                    State::SharedCta => {
48                        push_directive(tokens, "shared::cta");
49                    }
50                    State::Shared => {
51                        push_directive(tokens, "shared");
52                    }
53                }
54            }
55            push_directive(tokens, "b64");
56            self.state2.unparse_tokens(tokens);
57            tokens.push(PtxToken::Comma);
58            self.addr.unparse_tokens(tokens);
59            if self.count.is_some() {
60                tokens.push(PtxToken::Comma);
61            }
62            if let Some(opt_3) = self.count.as_ref() {
63                opt_3.unparse_tokens(tokens);
64            }
65            tokens.push(PtxToken::Semicolon);
66        }
67    }
68
69    impl PtxUnparser for MbarrierArriveDropSemScopeSharedClusterB64 {
70        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
71            push_opcode(tokens, "mbarrier");
72            push_directive(tokens, "arrive_drop");
73            if let Some(sem_4) = self.sem.as_ref() {
74                match sem_4 {
75                    Sem::Release => {
76                        push_directive(tokens, "release");
77                    }
78                    Sem::Relaxed => {
79                        push_directive(tokens, "relaxed");
80                    }
81                }
82            }
83            if let Some(scope_5) = self.scope.as_ref() {
84                match scope_5 {
85                    Scope::Cluster => {
86                        push_directive(tokens, "cluster");
87                    }
88                    Scope::Cta => {
89                        push_directive(tokens, "cta");
90                    }
91                }
92            }
93            if self.shared_cluster {
94                push_directive(tokens, "shared::cluster");
95            }
96            push_directive(tokens, "b64");
97            self.operand.unparse_tokens(tokens);
98            tokens.push(PtxToken::Comma);
99            self.addr.unparse_tokens(tokens);
100            if self.count.is_some() {
101                tokens.push(PtxToken::Comma);
102            }
103            if let Some(opt_6) = self.count.as_ref() {
104                opt_6.unparse_tokens(tokens);
105            }
106            tokens.push(PtxToken::Semicolon);
107        }
108    }
109
110    impl PtxUnparser for MbarrierArriveDropExpectTxStateSemScopeB64 {
111        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
112            push_opcode(tokens, "mbarrier");
113            push_directive(tokens, "arrive_drop");
114            push_directive(tokens, "expect_tx");
115            if let Some(state_7) = self.state.as_ref() {
116                match state_7 {
117                    State::SharedCta => {
118                        push_directive(tokens, "shared::cta");
119                    }
120                    State::Shared => {
121                        push_directive(tokens, "shared");
122                    }
123                }
124            }
125            if let Some(sem_8) = self.sem.as_ref() {
126                match sem_8 {
127                    Sem::Release => {
128                        push_directive(tokens, "release");
129                    }
130                    Sem::Relaxed => {
131                        push_directive(tokens, "relaxed");
132                    }
133                }
134            }
135            if let Some(scope_9) = self.scope.as_ref() {
136                match scope_9 {
137                    Scope::Cluster => {
138                        push_directive(tokens, "cluster");
139                    }
140                    Scope::Cta => {
141                        push_directive(tokens, "cta");
142                    }
143                }
144            }
145            push_directive(tokens, "b64");
146            self.state2.unparse_tokens(tokens);
147            tokens.push(PtxToken::Comma);
148            self.addr.unparse_tokens(tokens);
149            tokens.push(PtxToken::Comma);
150            self.tx_count.unparse_tokens(tokens);
151            tokens.push(PtxToken::Semicolon);
152        }
153    }
154
155    impl PtxUnparser for MbarrierArriveDropExpectTxSharedClusterSemScopeB64 {
156        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
157            push_opcode(tokens, "mbarrier");
158            push_directive(tokens, "arrive_drop");
159            push_directive(tokens, "expect_tx");
160            if self.shared_cluster {
161                push_directive(tokens, "shared::cluster");
162            }
163            if let Some(sem_10) = self.sem.as_ref() {
164                match sem_10 {
165                    Sem::Release => {
166                        push_directive(tokens, "release");
167                    }
168                    Sem::Relaxed => {
169                        push_directive(tokens, "relaxed");
170                    }
171                }
172            }
173            if let Some(scope_11) = self.scope.as_ref() {
174                match scope_11 {
175                    Scope::Cluster => {
176                        push_directive(tokens, "cluster");
177                    }
178                    Scope::Cta => {
179                        push_directive(tokens, "cta");
180                    }
181                }
182            }
183            push_directive(tokens, "b64");
184            self.operand.unparse_tokens(tokens);
185            tokens.push(PtxToken::Comma);
186            self.addr.unparse_tokens(tokens);
187            tokens.push(PtxToken::Comma);
188            self.tx_count.unparse_tokens(tokens);
189            tokens.push(PtxToken::Semicolon);
190        }
191    }
192
193    impl PtxUnparser for MbarrierArriveDropNocompleteReleaseCtaStateB64 {
194        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
195            push_opcode(tokens, "mbarrier");
196            push_directive(tokens, "arrive_drop");
197            push_directive(tokens, "noComplete");
198            if self.release {
199                push_directive(tokens, "release");
200            }
201            if self.cta {
202                push_directive(tokens, "cta");
203            }
204            if let Some(state_12) = self.state.as_ref() {
205                match state_12 {
206                    State::SharedCta => {
207                        push_directive(tokens, "shared::cta");
208                    }
209                    State::Shared => {
210                        push_directive(tokens, "shared");
211                    }
212                }
213            }
214            push_directive(tokens, "b64");
215            self.state2.unparse_tokens(tokens);
216            tokens.push(PtxToken::Comma);
217            self.addr.unparse_tokens(tokens);
218            tokens.push(PtxToken::Comma);
219            self.count.unparse_tokens(tokens);
220            tokens.push(PtxToken::Semicolon);
221        }
222    }
223}
ptx_parser/unparser/instruction/mbarrier_arrive_drop.rs

ptx_parser/unparser/instruction/
mbarrier_arrive_drop.rs
