ptx_parser/parser/instruction/
st_async.rs

1//! Original PTX specification:
2//!
3//! st.async{.sem}{.scope}{.ss}{.completion_mechanism}{.vec}.type [a], b, [mbar];
4//! .sem  =                 { .weak };
5//! .scope =                { .cluster };
6//! .ss   =                 { .shared::cluster };
7//! .type =                 { .b32, .b64,
8//! .u32, .u64,
9//! .s32, .s64,
10//! .f32, .f64 };
11//! .vec  =                 { .v2, .v4 };
12//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
13//! ---------------------------------------------------------
14//! st.async{.mmio}.sem.scope{.ss}.type [a], b;
15//! .sem =                  { .release };
16//! .scope =                { .gpu, .sys };
17//! .ss =                   { .global };
18//! .type =                 { .b8, .b16, .b32, .b64,
19//! .u8, .u16, .u32, .u64,
20//! .s8, .s16, .s32, .s64,
21//! .f32, .f64 };
22
23#![allow(unused)]
24
25use crate::lexer::PtxToken;
26use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
27use crate::r#type::common::*;
28
29pub mod section_0 {
30    use super::*;
31    use crate::r#type::instruction::st_async::section_0::*;
32
33    // ============================================================================
34    // Generated enum parsers
35    // ============================================================================
36
37    impl PtxParser for CompletionMechanism {
38        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
39            // Try MbarrierCompleteTxBytes
40            {
41                let saved_pos = stream.position();
42                if stream.expect_string(".mbarrier::complete_tx::bytes").is_ok() {
43                    return Ok(CompletionMechanism::MbarrierCompleteTxBytes);
44                }
45                stream.set_position(saved_pos);
46            }
47            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
48            let expected = &[".mbarrier::complete_tx::bytes"];
49            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
50            Err(crate::parser::unexpected_value(span, expected, found))
51        }
52    }
53
54    impl PtxParser for Scope {
55        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
56            // Try Cluster
57            {
58                let saved_pos = stream.position();
59                if stream.expect_string(".cluster").is_ok() {
60                    return Ok(Scope::Cluster);
61                }
62                stream.set_position(saved_pos);
63            }
64            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
65            let expected = &[".cluster"];
66            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
67            Err(crate::parser::unexpected_value(span, expected, found))
68        }
69    }
70
71    impl PtxParser for Sem {
72        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
73            // Try Weak
74            {
75                let saved_pos = stream.position();
76                if stream.expect_string(".weak").is_ok() {
77                    return Ok(Sem::Weak);
78                }
79                stream.set_position(saved_pos);
80            }
81            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
82            let expected = &[".weak"];
83            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
84            Err(crate::parser::unexpected_value(span, expected, found))
85        }
86    }
87
88    impl PtxParser for Ss {
89        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
90            // Try SharedCluster
91            {
92                let saved_pos = stream.position();
93                if stream.expect_string(".shared::cluster").is_ok() {
94                    return Ok(Ss::SharedCluster);
95                }
96                stream.set_position(saved_pos);
97            }
98            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
99            let expected = &[".shared::cluster"];
100            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
101            Err(crate::parser::unexpected_value(span, expected, found))
102        }
103    }
104
105    impl PtxParser for Type {
106        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
107            // Try B32
108            {
109                let saved_pos = stream.position();
110                if stream.expect_string(".b32").is_ok() {
111                    return Ok(Type::B32);
112                }
113                stream.set_position(saved_pos);
114            }
115            let saved_pos = stream.position();
116            // Try B64
117            {
118                let saved_pos = stream.position();
119                if stream.expect_string(".b64").is_ok() {
120                    return Ok(Type::B64);
121                }
122                stream.set_position(saved_pos);
123            }
124            stream.set_position(saved_pos);
125            let saved_pos = stream.position();
126            // Try U32
127            {
128                let saved_pos = stream.position();
129                if stream.expect_string(".u32").is_ok() {
130                    return Ok(Type::U32);
131                }
132                stream.set_position(saved_pos);
133            }
134            stream.set_position(saved_pos);
135            let saved_pos = stream.position();
136            // Try U64
137            {
138                let saved_pos = stream.position();
139                if stream.expect_string(".u64").is_ok() {
140                    return Ok(Type::U64);
141                }
142                stream.set_position(saved_pos);
143            }
144            stream.set_position(saved_pos);
145            let saved_pos = stream.position();
146            // Try S32
147            {
148                let saved_pos = stream.position();
149                if stream.expect_string(".s32").is_ok() {
150                    return Ok(Type::S32);
151                }
152                stream.set_position(saved_pos);
153            }
154            stream.set_position(saved_pos);
155            let saved_pos = stream.position();
156            // Try S64
157            {
158                let saved_pos = stream.position();
159                if stream.expect_string(".s64").is_ok() {
160                    return Ok(Type::S64);
161                }
162                stream.set_position(saved_pos);
163            }
164            stream.set_position(saved_pos);
165            let saved_pos = stream.position();
166            // Try F32
167            {
168                let saved_pos = stream.position();
169                if stream.expect_string(".f32").is_ok() {
170                    return Ok(Type::F32);
171                }
172                stream.set_position(saved_pos);
173            }
174            stream.set_position(saved_pos);
175            let saved_pos = stream.position();
176            // Try F64
177            {
178                let saved_pos = stream.position();
179                if stream.expect_string(".f64").is_ok() {
180                    return Ok(Type::F64);
181                }
182                stream.set_position(saved_pos);
183            }
184            stream.set_position(saved_pos);
185            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
186            let expected = &[".b32", ".b64", ".u32", ".u64", ".s32", ".s64", ".f32", ".f64"];
187            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
188            Err(crate::parser::unexpected_value(span, expected, found))
189        }
190    }
191
192    impl PtxParser for Vec {
193        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
194            // Try V2
195            {
196                let saved_pos = stream.position();
197                if stream.expect_string(".v2").is_ok() {
198                    return Ok(Vec::V2);
199                }
200                stream.set_position(saved_pos);
201            }
202            let saved_pos = stream.position();
203            // Try V4
204            {
205                let saved_pos = stream.position();
206                if stream.expect_string(".v4").is_ok() {
207                    return Ok(Vec::V4);
208                }
209                stream.set_position(saved_pos);
210            }
211            stream.set_position(saved_pos);
212            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
213            let expected = &[".v2", ".v4"];
214            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
215            Err(crate::parser::unexpected_value(span, expected, found))
216        }
217    }
218
219    impl PtxParser for StAsyncSemScopeSsCompletionMechanismVecType {
220        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
221            stream.expect_string("st")?;
222            stream.expect_string(".async")?;
223            let async_ = ();
224            stream.expect_complete()?;
225            let saved_pos = stream.position();
226            let sem = match Sem::parse(stream) {
227                Ok(val) => Some(val),
228                Err(_) => {
229                    stream.set_position(saved_pos);
230                    None
231                }
232            };
233            stream.expect_complete()?;
234            let saved_pos = stream.position();
235            let scope = match Scope::parse(stream) {
236                Ok(val) => Some(val),
237                Err(_) => {
238                    stream.set_position(saved_pos);
239                    None
240                }
241            };
242            stream.expect_complete()?;
243            let saved_pos = stream.position();
244            let ss = match Ss::parse(stream) {
245                Ok(val) => Some(val),
246                Err(_) => {
247                    stream.set_position(saved_pos);
248                    None
249                }
250            };
251            stream.expect_complete()?;
252            let saved_pos = stream.position();
253            let completion_mechanism = match CompletionMechanism::parse(stream) {
254                Ok(val) => Some(val),
255                Err(_) => {
256                    stream.set_position(saved_pos);
257                    None
258                }
259            };
260            stream.expect_complete()?;
261            let saved_pos = stream.position();
262            let vec = match Vec::parse(stream) {
263                Ok(val) => Some(val),
264                Err(_) => {
265                    stream.set_position(saved_pos);
266                    None
267                }
268            };
269            stream.expect_complete()?;
270            let type_ = Type::parse(stream)?;
271            stream.expect_complete()?;
272            let a = AddressOperand::parse(stream)?;
273            stream.expect_complete()?;
274            stream.expect(&PtxToken::Comma)?;
275            let b = GeneralOperand::parse(stream)?;
276            stream.expect_complete()?;
277            stream.expect(&PtxToken::Comma)?;
278            let mbar = AddressOperand::parse(stream)?;
279            stream.expect_complete()?;
280            stream.expect_complete()?;
281            stream.expect(&PtxToken::Semicolon)?;
282            Ok(StAsyncSemScopeSsCompletionMechanismVecType {
283                async_,
284                sem,
285                scope,
286                ss,
287                completion_mechanism,
288                vec,
289                type_,
290                a,
291                b,
292                mbar,
293            })
294        }
295    }
296
297
298}
299
300pub mod section_1 {
301    use super::*;
302    use crate::r#type::instruction::st_async::section_1::*;
303
304    // ============================================================================
305    // Generated enum parsers
306    // ============================================================================
307
308    impl PtxParser for Scope {
309        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
310            // Try Gpu
311            {
312                let saved_pos = stream.position();
313                if stream.expect_string(".gpu").is_ok() {
314                    return Ok(Scope::Gpu);
315                }
316                stream.set_position(saved_pos);
317            }
318            let saved_pos = stream.position();
319            // Try Sys
320            {
321                let saved_pos = stream.position();
322                if stream.expect_string(".sys").is_ok() {
323                    return Ok(Scope::Sys);
324                }
325                stream.set_position(saved_pos);
326            }
327            stream.set_position(saved_pos);
328            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
329            let expected = &[".gpu", ".sys"];
330            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
331            Err(crate::parser::unexpected_value(span, expected, found))
332        }
333    }
334
335    impl PtxParser for Sem {
336        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
337            // Try Release
338            {
339                let saved_pos = stream.position();
340                if stream.expect_string(".release").is_ok() {
341                    return Ok(Sem::Release);
342                }
343                stream.set_position(saved_pos);
344            }
345            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
346            let expected = &[".release"];
347            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
348            Err(crate::parser::unexpected_value(span, expected, found))
349        }
350    }
351
352    impl PtxParser for Ss {
353        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
354            // Try Global
355            {
356                let saved_pos = stream.position();
357                if stream.expect_string(".global").is_ok() {
358                    return Ok(Ss::Global);
359                }
360                stream.set_position(saved_pos);
361            }
362            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
363            let expected = &[".global"];
364            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
365            Err(crate::parser::unexpected_value(span, expected, found))
366        }
367    }
368
369    impl PtxParser for Type {
370        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
371            // Try B16
372            {
373                let saved_pos = stream.position();
374                if stream.expect_string(".b16").is_ok() {
375                    return Ok(Type::B16);
376                }
377                stream.set_position(saved_pos);
378            }
379            let saved_pos = stream.position();
380            // Try B32
381            {
382                let saved_pos = stream.position();
383                if stream.expect_string(".b32").is_ok() {
384                    return Ok(Type::B32);
385                }
386                stream.set_position(saved_pos);
387            }
388            stream.set_position(saved_pos);
389            let saved_pos = stream.position();
390            // Try B64
391            {
392                let saved_pos = stream.position();
393                if stream.expect_string(".b64").is_ok() {
394                    return Ok(Type::B64);
395                }
396                stream.set_position(saved_pos);
397            }
398            stream.set_position(saved_pos);
399            let saved_pos = stream.position();
400            // Try U16
401            {
402                let saved_pos = stream.position();
403                if stream.expect_string(".u16").is_ok() {
404                    return Ok(Type::U16);
405                }
406                stream.set_position(saved_pos);
407            }
408            stream.set_position(saved_pos);
409            let saved_pos = stream.position();
410            // Try U32
411            {
412                let saved_pos = stream.position();
413                if stream.expect_string(".u32").is_ok() {
414                    return Ok(Type::U32);
415                }
416                stream.set_position(saved_pos);
417            }
418            stream.set_position(saved_pos);
419            let saved_pos = stream.position();
420            // Try U64
421            {
422                let saved_pos = stream.position();
423                if stream.expect_string(".u64").is_ok() {
424                    return Ok(Type::U64);
425                }
426                stream.set_position(saved_pos);
427            }
428            stream.set_position(saved_pos);
429            let saved_pos = stream.position();
430            // Try S16
431            {
432                let saved_pos = stream.position();
433                if stream.expect_string(".s16").is_ok() {
434                    return Ok(Type::S16);
435                }
436                stream.set_position(saved_pos);
437            }
438            stream.set_position(saved_pos);
439            let saved_pos = stream.position();
440            // Try S32
441            {
442                let saved_pos = stream.position();
443                if stream.expect_string(".s32").is_ok() {
444                    return Ok(Type::S32);
445                }
446                stream.set_position(saved_pos);
447            }
448            stream.set_position(saved_pos);
449            let saved_pos = stream.position();
450            // Try S64
451            {
452                let saved_pos = stream.position();
453                if stream.expect_string(".s64").is_ok() {
454                    return Ok(Type::S64);
455                }
456                stream.set_position(saved_pos);
457            }
458            stream.set_position(saved_pos);
459            let saved_pos = stream.position();
460            // Try F32
461            {
462                let saved_pos = stream.position();
463                if stream.expect_string(".f32").is_ok() {
464                    return Ok(Type::F32);
465                }
466                stream.set_position(saved_pos);
467            }
468            stream.set_position(saved_pos);
469            let saved_pos = stream.position();
470            // Try F64
471            {
472                let saved_pos = stream.position();
473                if stream.expect_string(".f64").is_ok() {
474                    return Ok(Type::F64);
475                }
476                stream.set_position(saved_pos);
477            }
478            stream.set_position(saved_pos);
479            let saved_pos = stream.position();
480            // Try B8
481            {
482                let saved_pos = stream.position();
483                if stream.expect_string(".b8").is_ok() {
484                    return Ok(Type::B8);
485                }
486                stream.set_position(saved_pos);
487            }
488            stream.set_position(saved_pos);
489            let saved_pos = stream.position();
490            // Try U8
491            {
492                let saved_pos = stream.position();
493                if stream.expect_string(".u8").is_ok() {
494                    return Ok(Type::U8);
495                }
496                stream.set_position(saved_pos);
497            }
498            stream.set_position(saved_pos);
499            let saved_pos = stream.position();
500            // Try S8
501            {
502                let saved_pos = stream.position();
503                if stream.expect_string(".s8").is_ok() {
504                    return Ok(Type::S8);
505                }
506                stream.set_position(saved_pos);
507            }
508            stream.set_position(saved_pos);
509            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
510            let expected = &[".b16", ".b32", ".b64", ".u16", ".u32", ".u64", ".s16", ".s32", ".s64", ".f32", ".f64", ".b8", ".u8", ".s8"];
511            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
512            Err(crate::parser::unexpected_value(span, expected, found))
513        }
514    }
515
516    impl PtxParser for StAsyncMmioSemScopeSsType {
517        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
518            stream.expect_string("st")?;
519            stream.expect_string(".async")?;
520            let async_ = ();
521            stream.expect_complete()?;
522            let saved_pos = stream.position();
523            let mmio = stream.expect_string(".mmio").is_ok();
524            if !mmio {
525                stream.set_position(saved_pos);
526            }
527            stream.expect_complete()?;
528            let sem = Sem::parse(stream)?;
529            stream.expect_complete()?;
530            let scope = Scope::parse(stream)?;
531            stream.expect_complete()?;
532            let saved_pos = stream.position();
533            let ss = match Ss::parse(stream) {
534                Ok(val) => Some(val),
535                Err(_) => {
536                    stream.set_position(saved_pos);
537                    None
538                }
539            };
540            stream.expect_complete()?;
541            let type_ = Type::parse(stream)?;
542            stream.expect_complete()?;
543            let a = AddressOperand::parse(stream)?;
544            stream.expect_complete()?;
545            stream.expect(&PtxToken::Comma)?;
546            let b = GeneralOperand::parse(stream)?;
547            stream.expect_complete()?;
548            stream.expect_complete()?;
549            stream.expect(&PtxToken::Semicolon)?;
550            Ok(StAsyncMmioSemScopeSsType {
551                async_,
552                mmio,
553                sem,
554                scope,
555                ss,
556                type_,
557                a,
558                b,
559            })
560        }
561    }
562
563
564}
565
ptx_parser/parser/instruction/st_async.rs

ptx_parser/parser/instruction/
st_async.rs
