ptx_parser/parser/instruction/
tcgen05_st.rs

1//! Original PTX specification:
2//!
3//! tcgen05.st.sync.aligned.shape1.num{.unpack}.b32    [taddr], r;
4//! tcgen05.st.sync.aligned.shape2.num{.unpack}.b32    [taddr], immHalfSplitoff, r;
5//! .shape1 = { .16x64b, .16x128b, .16x256b, .32x32b };
6//! .shape2 = { .16x32bx2 };
7//! .num    = { .x1, .x2, .x4, .x8, .x16, .x32, .x64, .x128 };
8//! .unpack = { .unpack::16b };
9
10#![allow(unused)]
11
12use crate::lexer::PtxToken;
13use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
14use crate::r#type::common::*;
15
16pub mod section_0 {
17    use super::*;
18    use crate::r#type::instruction::tcgen05_st::section_0::*;
19
20    // ============================================================================
21    // Generated enum parsers
22    // ============================================================================
23
24    impl PtxParser for Num {
25        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
26            // Try X128
27            {
28                let saved_pos = stream.position();
29                if stream.expect_string(".x128").is_ok() {
30                    return Ok(Num::X128);
31                }
32                stream.set_position(saved_pos);
33            }
34            let saved_pos = stream.position();
35            // Try X16
36            {
37                let saved_pos = stream.position();
38                if stream.expect_string(".x16").is_ok() {
39                    return Ok(Num::X16);
40                }
41                stream.set_position(saved_pos);
42            }
43            stream.set_position(saved_pos);
44            let saved_pos = stream.position();
45            // Try X32
46            {
47                let saved_pos = stream.position();
48                if stream.expect_string(".x32").is_ok() {
49                    return Ok(Num::X32);
50                }
51                stream.set_position(saved_pos);
52            }
53            stream.set_position(saved_pos);
54            let saved_pos = stream.position();
55            // Try X64
56            {
57                let saved_pos = stream.position();
58                if stream.expect_string(".x64").is_ok() {
59                    return Ok(Num::X64);
60                }
61                stream.set_position(saved_pos);
62            }
63            stream.set_position(saved_pos);
64            let saved_pos = stream.position();
65            // Try X1
66            {
67                let saved_pos = stream.position();
68                if stream.expect_string(".x1").is_ok() {
69                    return Ok(Num::X1);
70                }
71                stream.set_position(saved_pos);
72            }
73            stream.set_position(saved_pos);
74            let saved_pos = stream.position();
75            // Try X2
76            {
77                let saved_pos = stream.position();
78                if stream.expect_string(".x2").is_ok() {
79                    return Ok(Num::X2);
80                }
81                stream.set_position(saved_pos);
82            }
83            stream.set_position(saved_pos);
84            let saved_pos = stream.position();
85            // Try X4
86            {
87                let saved_pos = stream.position();
88                if stream.expect_string(".x4").is_ok() {
89                    return Ok(Num::X4);
90                }
91                stream.set_position(saved_pos);
92            }
93            stream.set_position(saved_pos);
94            let saved_pos = stream.position();
95            // Try X8
96            {
97                let saved_pos = stream.position();
98                if stream.expect_string(".x8").is_ok() {
99                    return Ok(Num::X8);
100                }
101                stream.set_position(saved_pos);
102            }
103            stream.set_position(saved_pos);
104            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
105            let expected = &[".x128", ".x16", ".x32", ".x64", ".x1", ".x2", ".x4", ".x8"];
106            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
107            Err(crate::parser::unexpected_value(span, expected, found))
108        }
109    }
110
111    impl PtxParser for Shape1 {
112        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
113            // Try _16x128b
114            {
115                let saved_pos = stream.position();
116                if stream.expect_string(".16x128b").is_ok() {
117                    return Ok(Shape1::_16x128b);
118                }
119                stream.set_position(saved_pos);
120            }
121            let saved_pos = stream.position();
122            // Try _16x256b
123            {
124                let saved_pos = stream.position();
125                if stream.expect_string(".16x256b").is_ok() {
126                    return Ok(Shape1::_16x256b);
127                }
128                stream.set_position(saved_pos);
129            }
130            stream.set_position(saved_pos);
131            let saved_pos = stream.position();
132            // Try _16x64b
133            {
134                let saved_pos = stream.position();
135                if stream.expect_string(".16x64b").is_ok() {
136                    return Ok(Shape1::_16x64b);
137                }
138                stream.set_position(saved_pos);
139            }
140            stream.set_position(saved_pos);
141            let saved_pos = stream.position();
142            // Try _32x32b
143            {
144                let saved_pos = stream.position();
145                if stream.expect_string(".32x32b").is_ok() {
146                    return Ok(Shape1::_32x32b);
147                }
148                stream.set_position(saved_pos);
149            }
150            stream.set_position(saved_pos);
151            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
152            let expected = &[".16x128b", ".16x256b", ".16x64b", ".32x32b"];
153            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
154            Err(crate::parser::unexpected_value(span, expected, found))
155        }
156    }
157
158    impl PtxParser for Shape2 {
159        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
160            // Try _16x32bx2
161            {
162                let saved_pos = stream.position();
163                if stream.expect_string(".16x32bx2").is_ok() {
164                    return Ok(Shape2::_16x32bx2);
165                }
166                stream.set_position(saved_pos);
167            }
168            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
169            let expected = &[".16x32bx2"];
170            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
171            Err(crate::parser::unexpected_value(span, expected, found))
172        }
173    }
174
175    impl PtxParser for Unpack {
176        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
177            // Try Unpack16b
178            {
179                let saved_pos = stream.position();
180                if stream.expect_string(".unpack::16b").is_ok() {
181                    return Ok(Unpack::Unpack16b);
182                }
183                stream.set_position(saved_pos);
184            }
185            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
186            let expected = &[".unpack::16b"];
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 Tcgen05StSyncAlignedShape1NumUnpackB32 {
193        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
194            stream.expect_string("tcgen05")?;
195            stream.expect_string(".st")?;
196            let st = ();
197            stream.expect_complete()?;
198            stream.expect_string(".sync")?;
199            let sync = ();
200            stream.expect_complete()?;
201            stream.expect_string(".aligned")?;
202            let aligned = ();
203            stream.expect_complete()?;
204            let shape1 = Shape1::parse(stream)?;
205            stream.expect_complete()?;
206            let num = Num::parse(stream)?;
207            stream.expect_complete()?;
208            let saved_pos = stream.position();
209            let unpack = match Unpack::parse(stream) {
210                Ok(val) => Some(val),
211                Err(_) => {
212                    stream.set_position(saved_pos);
213                    None
214                }
215            };
216            stream.expect_complete()?;
217            stream.expect_string(".b32")?;
218            let b32 = ();
219            stream.expect_complete()?;
220            let taddr = AddressOperand::parse(stream)?;
221            stream.expect_complete()?;
222            stream.expect(&PtxToken::Comma)?;
223            let r = GeneralOperand::parse(stream)?;
224            stream.expect_complete()?;
225            stream.expect_complete()?;
226            stream.expect(&PtxToken::Semicolon)?;
227            Ok(Tcgen05StSyncAlignedShape1NumUnpackB32 {
228                st,
229                sync,
230                aligned,
231                shape1,
232                num,
233                unpack,
234                b32,
235                taddr,
236                r,
237            })
238        }
239    }
240
241
242    impl PtxParser for Tcgen05StSyncAlignedShape2NumUnpackB32 {
243        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
244            stream.expect_string("tcgen05")?;
245            stream.expect_string(".st")?;
246            let st = ();
247            stream.expect_complete()?;
248            stream.expect_string(".sync")?;
249            let sync = ();
250            stream.expect_complete()?;
251            stream.expect_string(".aligned")?;
252            let aligned = ();
253            stream.expect_complete()?;
254            let shape2 = Shape2::parse(stream)?;
255            stream.expect_complete()?;
256            let num = Num::parse(stream)?;
257            stream.expect_complete()?;
258            let saved_pos = stream.position();
259            let unpack = match Unpack::parse(stream) {
260                Ok(val) => Some(val),
261                Err(_) => {
262                    stream.set_position(saved_pos);
263                    None
264                }
265            };
266            stream.expect_complete()?;
267            stream.expect_string(".b32")?;
268            let b32 = ();
269            stream.expect_complete()?;
270            let taddr = AddressOperand::parse(stream)?;
271            stream.expect_complete()?;
272            stream.expect(&PtxToken::Comma)?;
273            let immhalfsplitoff = GeneralOperand::parse(stream)?;
274            stream.expect_complete()?;
275            stream.expect(&PtxToken::Comma)?;
276            let r = GeneralOperand::parse(stream)?;
277            stream.expect_complete()?;
278            stream.expect_complete()?;
279            stream.expect(&PtxToken::Semicolon)?;
280            Ok(Tcgen05StSyncAlignedShape2NumUnpackB32 {
281                st,
282                sync,
283                aligned,
284                shape2,
285                num,
286                unpack,
287                b32,
288                taddr,
289                immhalfsplitoff,
290                r,
291            })
292        }
293    }
294
295
296}
297
ptx_parser/parser/instruction/tcgen05_st.rs

ptx_parser/parser/instruction/
tcgen05_st.rs
