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.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
43            let expected = &[".cta_group::1", ".cta_group::2"];
44            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
45            Err(crate::parser::unexpected_value(span, expected, found))
46        }
47    }
48
49    impl PtxParser for Tcgen05AllocCtaGroupSyncAlignedSharedCtaB32 {
50        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
51            stream.expect_string("tcgen05")?;
52            stream.expect_string(".alloc")?;
53            let alloc = ();
54            stream.expect_complete()?;
55            let cta_group = CtaGroup::parse(stream)?;
56            stream.expect_complete()?;
57            stream.expect_string(".sync")?;
58            let sync = ();
59            stream.expect_complete()?;
60            stream.expect_string(".aligned")?;
61            let aligned = ();
62            stream.expect_complete()?;
63            let saved_pos = stream.position();
64            let shared_cta = stream.expect_string(".shared::cta").is_ok();
65            if !shared_cta {
66                stream.set_position(saved_pos);
67            }
68            stream.expect_complete()?;
69            stream.expect_string(".b32")?;
70            let b32 = ();
71            stream.expect_complete()?;
72            let dst = AddressOperand::parse(stream)?;
73            stream.expect_complete()?;
74            stream.expect(&PtxToken::Comma)?;
75            let ncols = GeneralOperand::parse(stream)?;
76            stream.expect_complete()?;
77            stream.expect_complete()?;
78            stream.expect(&PtxToken::Semicolon)?;
79            Ok(Tcgen05AllocCtaGroupSyncAlignedSharedCtaB32 {
80                alloc,
81                cta_group,
82                sync,
83                aligned,
84                shared_cta,
85                b32,
86                dst,
87                ncols,
88            })
89        }
90    }
91
92
93    impl PtxParser for Tcgen05DeallocCtaGroupSyncAlignedB32 {
94        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
95            stream.expect_string("tcgen05")?;
96            stream.expect_string(".dealloc")?;
97            let dealloc = ();
98            stream.expect_complete()?;
99            let cta_group = CtaGroup::parse(stream)?;
100            stream.expect_complete()?;
101            stream.expect_string(".sync")?;
102            let sync = ();
103            stream.expect_complete()?;
104            stream.expect_string(".aligned")?;
105            let aligned = ();
106            stream.expect_complete()?;
107            stream.expect_string(".b32")?;
108            let b32 = ();
109            stream.expect_complete()?;
110            let taddr = GeneralOperand::parse(stream)?;
111            stream.expect_complete()?;
112            stream.expect(&PtxToken::Comma)?;
113            let ncols = GeneralOperand::parse(stream)?;
114            stream.expect_complete()?;
115            stream.expect_complete()?;
116            stream.expect(&PtxToken::Semicolon)?;
117            Ok(Tcgen05DeallocCtaGroupSyncAlignedB32 {
118                dealloc,
119                cta_group,
120                sync,
121                aligned,
122                b32,
123                taddr,
124                ncols,
125            })
126        }
127    }
128
129
130    impl PtxParser for Tcgen05RelinquishAllocPermitCtaGroupSyncAligned {
131        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
132            stream.expect_string("tcgen05")?;
133            stream.expect_string(".relinquish_alloc_permit")?;
134            let relinquish_alloc_permit = ();
135            stream.expect_complete()?;
136            let cta_group = CtaGroup::parse(stream)?;
137            stream.expect_complete()?;
138            stream.expect_string(".sync")?;
139            let sync = ();
140            stream.expect_complete()?;
141            stream.expect_string(".aligned")?;
142            let aligned = ();
143            stream.expect_complete()?;
144            stream.expect_complete()?;
145            stream.expect(&PtxToken::Semicolon)?;
146            Ok(Tcgen05RelinquishAllocPermitCtaGroupSyncAligned {
147                relinquish_alloc_permit,
148                cta_group,
149                sync,
150                aligned,
151            })
152        }
153    }
154
155
156}
157
ptx_parser/parser/instruction/tcgen05_alloc.rs

ptx_parser/parser/instruction/
tcgen05_alloc.rs
