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            push_opcode(tokens, "mbarrier");
23            push_directive(tokens, "test_wait");
24            if let Some(sem_0) = self.sem.as_ref() {
25                match sem_0 {
26                    Sem::Acquire => {
27                        push_directive(tokens, "acquire");
28                    }
29                    Sem::Relaxed => {
30                        push_directive(tokens, "relaxed");
31                    }
32                }
33            }
34            if let Some(scope_1) = self.scope.as_ref() {
35                match scope_1 {
36                    Scope::Cluster => {
37                        push_directive(tokens, "cluster");
38                    }
39                    Scope::Cta => {
40                        push_directive(tokens, "cta");
41                    }
42                }
43            }
44            if let Some(state_2) = self.state.as_ref() {
45                match state_2 {
46                    State::SharedCta => {
47                        push_directive(tokens, "shared::cta");
48                    }
49                    State::Shared => {
50                        push_directive(tokens, "shared");
51                    }
52                }
53            }
54            push_directive(tokens, "b64");
55            self.waitcomplete.unparse_tokens(tokens);
56            tokens.push(PtxToken::Comma);
57            self.addr.unparse_tokens(tokens);
58            tokens.push(PtxToken::Comma);
59            self.state2.unparse_tokens(tokens);
60            tokens.push(PtxToken::Semicolon);
61        }
62    }
63
64    impl PtxUnparser for MbarrierTestWaitParitySemScopeStateB64 {
65        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
66            push_opcode(tokens, "mbarrier");
67            push_directive(tokens, "test_wait");
68            push_directive(tokens, "parity");
69            if let Some(sem_3) = self.sem.as_ref() {
70                match sem_3 {
71                    Sem::Acquire => {
72                        push_directive(tokens, "acquire");
73                    }
74                    Sem::Relaxed => {
75                        push_directive(tokens, "relaxed");
76                    }
77                }
78            }
79            if let Some(scope_4) = self.scope.as_ref() {
80                match scope_4 {
81                    Scope::Cluster => {
82                        push_directive(tokens, "cluster");
83                    }
84                    Scope::Cta => {
85                        push_directive(tokens, "cta");
86                    }
87                }
88            }
89            if let Some(state_5) = self.state.as_ref() {
90                match state_5 {
91                    State::SharedCta => {
92                        push_directive(tokens, "shared::cta");
93                    }
94                    State::Shared => {
95                        push_directive(tokens, "shared");
96                    }
97                }
98            }
99            push_directive(tokens, "b64");
100            self.waitcomplete.unparse_tokens(tokens);
101            tokens.push(PtxToken::Comma);
102            self.addr.unparse_tokens(tokens);
103            tokens.push(PtxToken::Comma);
104            self.phaseparity.unparse_tokens(tokens);
105            tokens.push(PtxToken::Semicolon);
106        }
107    }
108
109    impl PtxUnparser for MbarrierTryWaitSemScopeStateB64 {
110        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
111            push_opcode(tokens, "mbarrier");
112            push_directive(tokens, "try_wait");
113            if let Some(sem_6) = self.sem.as_ref() {
114                match sem_6 {
115                    Sem::Acquire => {
116                        push_directive(tokens, "acquire");
117                    }
118                    Sem::Relaxed => {
119                        push_directive(tokens, "relaxed");
120                    }
121                }
122            }
123            if let Some(scope_7) = self.scope.as_ref() {
124                match scope_7 {
125                    Scope::Cluster => {
126                        push_directive(tokens, "cluster");
127                    }
128                    Scope::Cta => {
129                        push_directive(tokens, "cta");
130                    }
131                }
132            }
133            if let Some(state_8) = self.state.as_ref() {
134                match state_8 {
135                    State::SharedCta => {
136                        push_directive(tokens, "shared::cta");
137                    }
138                    State::Shared => {
139                        push_directive(tokens, "shared");
140                    }
141                }
142            }
143            push_directive(tokens, "b64");
144            self.waitcomplete.unparse_tokens(tokens);
145            tokens.push(PtxToken::Comma);
146            self.addr.unparse_tokens(tokens);
147            tokens.push(PtxToken::Comma);
148            self.state2.unparse_tokens(tokens);
149            if self.suspendtimehint.is_some() {
150                tokens.push(PtxToken::Comma);
151            }
152            if let Some(opt_9) = self.suspendtimehint.as_ref() {
153                opt_9.unparse_tokens(tokens);
154            }
155            tokens.push(PtxToken::Semicolon);
156        }
157    }
158
159    impl PtxUnparser for MbarrierTryWaitParitySemScopeStateB64 {
160        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
161            push_opcode(tokens, "mbarrier");
162            push_directive(tokens, "try_wait");
163            push_directive(tokens, "parity");
164            if let Some(sem_10) = self.sem.as_ref() {
165                match sem_10 {
166                    Sem::Acquire => {
167                        push_directive(tokens, "acquire");
168                    }
169                    Sem::Relaxed => {
170                        push_directive(tokens, "relaxed");
171                    }
172                }
173            }
174            if let Some(scope_11) = self.scope.as_ref() {
175                match scope_11 {
176                    Scope::Cluster => {
177                        push_directive(tokens, "cluster");
178                    }
179                    Scope::Cta => {
180                        push_directive(tokens, "cta");
181                    }
182                }
183            }
184            if let Some(state_12) = self.state.as_ref() {
185                match state_12 {
186                    State::SharedCta => {
187                        push_directive(tokens, "shared::cta");
188                    }
189                    State::Shared => {
190                        push_directive(tokens, "shared");
191                    }
192                }
193            }
194            push_directive(tokens, "b64");
195            self.waitcomplete.unparse_tokens(tokens);
196            tokens.push(PtxToken::Comma);
197            self.addr.unparse_tokens(tokens);
198            tokens.push(PtxToken::Comma);
199            self.phaseparity.unparse_tokens(tokens);
200            if self.suspendtimehint.is_some() {
201                tokens.push(PtxToken::Comma);
202            }
203            if let Some(opt_13) = self.suspendtimehint.as_ref() {
204                opt_13.unparse_tokens(tokens);
205            }
206            tokens.push(PtxToken::Semicolon);
207        }
208    }
209}
ptx_parser/unparser/instruction/mbarrier_test_wait.rs

ptx_parser/unparser/instruction/
mbarrier_test_wait.rs
