ptx_parser/parser/instruction/
tcgen05_alloc.rs

1//! Original PTX specification:
2//!
3//! tcgen05.alloc.cta_group.sync.aligned{.shared::cta}.b32  [dst], nCols;
4//! tcgen05.dealloc.cta_group.sync.aligned.b32              taddr, nCols;
5//! tcgen05.relinquish_alloc_permit.cta_group.sync.aligned;
6//! .cta_group = { .cta_group::1, .cta_group::2 };
7
8#![allow(unused)]
9
10use crate::lexer::PtxToken;
11use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
12use crate::r#type::common::*;
13
14pub mod section_0 {
15    use super::*;
16    use crate::r#type::instruction::tcgen05_alloc::section_0::*;
17
18    // ============================================================================
19    // Generated enum parsers
20    // ============================================================================
21
22    impl PtxParser for CtaGroup {
23        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
24            // Try CtaGroup1
25            {
26                let saved_pos = stream.position();
27                if stream.expect_string(".cta_group::1").is_ok() {
28                    return Ok(CtaGroup::CtaGroup1);
29                }
30                stream.set_position(saved_pos);
31            }
32            let saved_pos = stream.position();
33            // Try CtaGroup2
34            {
35                let saved_pos = stream.position();
36                if stream.expect_string(".cta_group::2").is_ok() {
37                    return Ok(CtaGroup::CtaGroup2);
38                }
39                stream.set_position(saved_pos);
40            }
41            stream.set_position(saved_pos);
42            let span = stream
43                .peek()
44                .map(|(_, s)| s.clone())
45                .unwrap_or(Span { start: 0, end: 0 });
46            let expected = &[".cta_group::1", ".cta_group::2"];
47            let found = stream
48                .peek()
49                .map(|(t, _)| format!("{:?}", t))
50                .unwrap_or_else(|_| "<end of input>".to_string());
51            Err(crate::parser::unexpected_value(span, expected, found))
52        }
53    }
54
55    impl PtxParser for Tcgen05AllocCtaGroupSyncAlignedSharedCtaB32 {
56        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
57            stream.expect_string("tcgen05")?;
58            stream.expect_string(".alloc")?;
59            let alloc = ();
60            stream.expect_complete()?;
61            let cta_group = CtaGroup::parse(stream)?;
62            stream.expect_complete()?;
63            stream.expect_string(".sync")?;
64            let sync = ();
65            stream.expect_complete()?;
66            stream.expect_string(".aligned")?;
67            let aligned = ();
68            stream.expect_complete()?;
69            let saved_pos = stream.position();
70            let shared_cta = stream.expect_string(".shared::cta").is_ok();
71            if !shared_cta {
72                stream.set_position(saved_pos);
73            }
74            stream.expect_complete()?;
75            stream.expect_string(".b32")?;
76            let b32 = ();
77            stream.expect_complete()?;
78            let dst = AddressOperand::parse(stream)?;
79            stream.expect_complete()?;
80            stream.expect(&PtxToken::Comma)?;
81            let ncols = GeneralOperand::parse(stream)?;
82            stream.expect_complete()?;
83            stream.expect_complete()?;
84            stream.expect(&PtxToken::Semicolon)?;
85            Ok(Tcgen05AllocCtaGroupSyncAlignedSharedCtaB32 {
86                alloc,
87                cta_group,
88                sync,
89                aligned,
90                shared_cta,
91                b32,
92                dst,
93                ncols,
94            })
95        }
96    }
97
98    impl PtxParser for Tcgen05DeallocCtaGroupSyncAlignedB32 {
99        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
100            stream.expect_string("tcgen05")?;
101            stream.expect_string(".dealloc")?;
102            let dealloc = ();
103            stream.expect_complete()?;
104            let cta_group = CtaGroup::parse(stream)?;
105            stream.expect_complete()?;
106            stream.expect_string(".sync")?;
107            let sync = ();
108            stream.expect_complete()?;
109            stream.expect_string(".aligned")?;
110            let aligned = ();
111            stream.expect_complete()?;
112            stream.expect_string(".b32")?;
113            let b32 = ();
114            stream.expect_complete()?;
115            let taddr = GeneralOperand::parse(stream)?;
116            stream.expect_complete()?;
117            stream.expect(&PtxToken::Comma)?;
118            let ncols = GeneralOperand::parse(stream)?;
119            stream.expect_complete()?;
120            stream.expect_complete()?;
121            stream.expect(&PtxToken::Semicolon)?;
122            Ok(Tcgen05DeallocCtaGroupSyncAlignedB32 {
123                dealloc,
124                cta_group,
125                sync,
126                aligned,
127                b32,
128                taddr,
129                ncols,
130            })
131        }
132    }
133
134    impl PtxParser for Tcgen05RelinquishAllocPermitCtaGroupSyncAligned {
135        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
136            stream.expect_string("tcgen05")?;
137            stream.expect_string(".relinquish_alloc_permit")?;
138            let relinquish_alloc_permit = ();
139            stream.expect_complete()?;
140            let cta_group = CtaGroup::parse(stream)?;
141            stream.expect_complete()?;
142            stream.expect_string(".sync")?;
143            let sync = ();
144            stream.expect_complete()?;
145            stream.expect_string(".aligned")?;
146            let aligned = ();
147            stream.expect_complete()?;
148            stream.expect_complete()?;
149            stream.expect(&PtxToken::Semicolon)?;
150            Ok(Tcgen05RelinquishAllocPermitCtaGroupSyncAligned {
151                relinquish_alloc_permit,
152                cta_group,
153                sync,
154                aligned,
155            })
156        }
157    }
158}
ptx_parser/parser/instruction/tcgen05_alloc.rs

ptx_parser/parser/instruction/
tcgen05_alloc.rs
