ptx_parser/parser/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::parser::{
14    PtxParseError, PtxParser, PtxTokenStream, Span,
15    util::{
16        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
17        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
18    },
19};
20use crate::r#type::common::*;
21use crate::{alt, ok, seq_n};
22
23pub mod section_0 {
24    use super::*;
25    use crate::r#type::instruction::mbarrier_test_wait::section_0::*;
26
27    // ============================================================================
28    // Generated enum parsers
29    // ============================================================================
30
31    impl PtxParser for Scope {
32        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
33            alt!(
34                map(string_p(".cluster"), |_, _span| Scope::Cluster),
35                map(string_p(".cta"), |_, _span| Scope::Cta)
36            )
37        }
38    }
39
40    impl PtxParser for Sem {
41        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
42            alt!(
43                map(string_p(".acquire"), |_, _span| Sem::Acquire),
44                map(string_p(".relaxed"), |_, _span| Sem::Relaxed)
45            )
46        }
47    }
48
49    impl PtxParser for State {
50        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
51            alt!(
52                map(string_p(".shared::cta"), |_, _span| State::SharedCta),
53                map(string_p(".shared"), |_, _span| State::Shared)
54            )
55        }
56    }
57
58    impl PtxParser for MbarrierTestWaitSemScopeStateB64 {
59        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
60            try_map(
61                seq_n!(
62                    string_p("mbarrier"),
63                    string_p(".test_wait"),
64                    optional(Sem::parse()),
65                    optional(Scope::parse()),
66                    optional(State::parse()),
67                    string_p(".b64"),
68                    GeneralOperand::parse(),
69                    comma_p(),
70                    AddressOperand::parse(),
71                    comma_p(),
72                    GeneralOperand::parse(),
73                    semicolon_p()
74                ),
75                |(_, test_wait, sem, scope, state, b64, waitcomplete, _, addr, _, state2, _),
76                 span| {
77                    ok!(MbarrierTestWaitSemScopeStateB64 {
78                        test_wait = test_wait,
79                        sem = sem,
80                        scope = scope,
81                        state = state,
82                        b64 = b64,
83                        waitcomplete = waitcomplete,
84                        addr = addr,
85                        state2 = state2,
86
87                    })
88                },
89            )
90        }
91    }
92
93    impl PtxParser for MbarrierTestWaitParitySemScopeStateB64 {
94        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
95            try_map(
96                seq_n!(
97                    string_p("mbarrier"),
98                    string_p(".test_wait"),
99                    string_p(".parity"),
100                    optional(Sem::parse()),
101                    optional(Scope::parse()),
102                    optional(State::parse()),
103                    string_p(".b64"),
104                    GeneralOperand::parse(),
105                    comma_p(),
106                    AddressOperand::parse(),
107                    comma_p(),
108                    GeneralOperand::parse(),
109                    semicolon_p()
110                ),
111                |(
112                    _,
113                    test_wait,
114                    parity,
115                    sem,
116                    scope,
117                    state,
118                    b64,
119                    waitcomplete,
120                    _,
121                    addr,
122                    _,
123                    phaseparity,
124                    _,
125                ),
126                 span| {
127                    ok!(MbarrierTestWaitParitySemScopeStateB64 {
128                        test_wait = test_wait,
129                        parity = parity,
130                        sem = sem,
131                        scope = scope,
132                        state = state,
133                        b64 = b64,
134                        waitcomplete = waitcomplete,
135                        addr = addr,
136                        phaseparity = phaseparity,
137
138                    })
139                },
140            )
141        }
142    }
143
144    impl PtxParser for MbarrierTryWaitSemScopeStateB64 {
145        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
146            try_map(
147                seq_n!(
148                    string_p("mbarrier"),
149                    string_p(".try_wait"),
150                    optional(Sem::parse()),
151                    optional(Scope::parse()),
152                    optional(State::parse()),
153                    string_p(".b64"),
154                    GeneralOperand::parse(),
155                    comma_p(),
156                    AddressOperand::parse(),
157                    comma_p(),
158                    GeneralOperand::parse(),
159                    map(
160                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
161                        |value, _| value.map(|(_, operand)| operand)
162                    ),
163                    semicolon_p()
164                ),
165                |(
166                    _,
167                    try_wait,
168                    sem,
169                    scope,
170                    state,
171                    b64,
172                    waitcomplete,
173                    _,
174                    addr,
175                    _,
176                    state2,
177                    suspendtimehint,
178                    _,
179                ),
180                 span| {
181                    ok!(MbarrierTryWaitSemScopeStateB64 {
182                        try_wait = try_wait,
183                        sem = sem,
184                        scope = scope,
185                        state = state,
186                        b64 = b64,
187                        waitcomplete = waitcomplete,
188                        addr = addr,
189                        state2 = state2,
190                        suspendtimehint = suspendtimehint,
191
192                    })
193                },
194            )
195        }
196    }
197
198    impl PtxParser for MbarrierTryWaitParitySemScopeStateB64 {
199        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
200            try_map(
201                seq_n!(
202                    string_p("mbarrier"),
203                    string_p(".try_wait"),
204                    string_p(".parity"),
205                    optional(Sem::parse()),
206                    optional(Scope::parse()),
207                    optional(State::parse()),
208                    string_p(".b64"),
209                    GeneralOperand::parse(),
210                    comma_p(),
211                    AddressOperand::parse(),
212                    comma_p(),
213                    GeneralOperand::parse(),
214                    map(
215                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
216                        |value, _| value.map(|(_, operand)| operand)
217                    ),
218                    semicolon_p()
219                ),
220                |(
221                    _,
222                    try_wait,
223                    parity,
224                    sem,
225                    scope,
226                    state,
227                    b64,
228                    waitcomplete,
229                    _,
230                    addr,
231                    _,
232                    phaseparity,
233                    suspendtimehint,
234                    _,
235                ),
236                 span| {
237                    ok!(MbarrierTryWaitParitySemScopeStateB64 {
238                        try_wait = try_wait,
239                        parity = parity,
240                        sem = sem,
241                        scope = scope,
242                        state = state,
243                        b64 = b64,
244                        waitcomplete = waitcomplete,
245                        addr = addr,
246                        phaseparity = phaseparity,
247                        suspendtimehint = suspendtimehint,
248
249                    })
250                },
251            )
252        }
253    }
254}
ptx_parser/parser/instruction/mbarrier_test_wait.rs

ptx_parser/parser/instruction/
mbarrier_test_wait.rs
