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

ptx_parser/parser/instruction/
st_async.rs
