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
105                .peek()
106                .map(|(_, s)| s.clone())
107                .unwrap_or(Span { start: 0, end: 0 });
108            let expected = &[".x128", ".x16", ".x32", ".x64", ".x1", ".x2", ".x4", ".x8"];
109            let found = stream
110                .peek()
111                .map(|(t, _)| format!("{:?}", t))
112                .unwrap_or_else(|_| "<end of input>".to_string());
113            Err(crate::parser::unexpected_value(span, expected, found))
114        }
115    }
116
117    impl PtxParser for Shape1 {
118        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
119            // Try _16x128b
120            {
121                let saved_pos = stream.position();
122                if stream.expect_string(".16x128b").is_ok() {
123                    return Ok(Shape1::_16x128b);
124                }
125                stream.set_position(saved_pos);
126            }
127            let saved_pos = stream.position();
128            // Try _16x256b
129            {
130                let saved_pos = stream.position();
131                if stream.expect_string(".16x256b").is_ok() {
132                    return Ok(Shape1::_16x256b);
133                }
134                stream.set_position(saved_pos);
135            }
136            stream.set_position(saved_pos);
137            let saved_pos = stream.position();
138            // Try _16x64b
139            {
140                let saved_pos = stream.position();
141                if stream.expect_string(".16x64b").is_ok() {
142                    return Ok(Shape1::_16x64b);
143                }
144                stream.set_position(saved_pos);
145            }
146            stream.set_position(saved_pos);
147            let saved_pos = stream.position();
148            // Try _32x32b
149            {
150                let saved_pos = stream.position();
151                if stream.expect_string(".32x32b").is_ok() {
152                    return Ok(Shape1::_32x32b);
153                }
154                stream.set_position(saved_pos);
155            }
156            stream.set_position(saved_pos);
157            let span = stream
158                .peek()
159                .map(|(_, s)| s.clone())
160                .unwrap_or(Span { start: 0, end: 0 });
161            let expected = &[".16x128b", ".16x256b", ".16x64b", ".32x32b"];
162            let found = stream
163                .peek()
164                .map(|(t, _)| format!("{:?}", t))
165                .unwrap_or_else(|_| "<end of input>".to_string());
166            Err(crate::parser::unexpected_value(span, expected, found))
167        }
168    }
169
170    impl PtxParser for Shape2 {
171        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
172            // Try _16x32bx2
173            {
174                let saved_pos = stream.position();
175                if stream.expect_string(".16x32bx2").is_ok() {
176                    return Ok(Shape2::_16x32bx2);
177                }
178                stream.set_position(saved_pos);
179            }
180            let span = stream
181                .peek()
182                .map(|(_, s)| s.clone())
183                .unwrap_or(Span { start: 0, end: 0 });
184            let expected = &[".16x32bx2"];
185            let found = stream
186                .peek()
187                .map(|(t, _)| format!("{:?}", t))
188                .unwrap_or_else(|_| "<end of input>".to_string());
189            Err(crate::parser::unexpected_value(span, expected, found))
190        }
191    }
192
193    impl PtxParser for Unpack {
194        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
195            // Try Unpack16b
196            {
197                let saved_pos = stream.position();
198                if stream.expect_string(".unpack::16b").is_ok() {
199                    return Ok(Unpack::Unpack16b);
200                }
201                stream.set_position(saved_pos);
202            }
203            let span = stream
204                .peek()
205                .map(|(_, s)| s.clone())
206                .unwrap_or(Span { start: 0, end: 0 });
207            let expected = &[".unpack::16b"];
208            let found = stream
209                .peek()
210                .map(|(t, _)| format!("{:?}", t))
211                .unwrap_or_else(|_| "<end of input>".to_string());
212            Err(crate::parser::unexpected_value(span, expected, found))
213        }
214    }
215
216    impl PtxParser for Tcgen05StSyncAlignedShape1NumUnpackB32 {
217        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
218            stream.expect_string("tcgen05")?;
219            stream.expect_string(".st")?;
220            let st = ();
221            stream.expect_complete()?;
222            stream.expect_string(".sync")?;
223            let sync = ();
224            stream.expect_complete()?;
225            stream.expect_string(".aligned")?;
226            let aligned = ();
227            stream.expect_complete()?;
228            let shape1 = Shape1::parse(stream)?;
229            stream.expect_complete()?;
230            let num = Num::parse(stream)?;
231            stream.expect_complete()?;
232            let saved_pos = stream.position();
233            let unpack = match Unpack::parse(stream) {
234                Ok(val) => Some(val),
235                Err(_) => {
236                    stream.set_position(saved_pos);
237                    None
238                }
239            };
240            stream.expect_complete()?;
241            stream.expect_string(".b32")?;
242            let b32 = ();
243            stream.expect_complete()?;
244            let taddr = AddressOperand::parse(stream)?;
245            stream.expect_complete()?;
246            stream.expect(&PtxToken::Comma)?;
247            let r = GeneralOperand::parse(stream)?;
248            stream.expect_complete()?;
249            stream.expect_complete()?;
250            stream.expect(&PtxToken::Semicolon)?;
251            Ok(Tcgen05StSyncAlignedShape1NumUnpackB32 {
252                st,
253                sync,
254                aligned,
255                shape1,
256                num,
257                unpack,
258                b32,
259                taddr,
260                r,
261            })
262        }
263    }
264
265    impl PtxParser for Tcgen05StSyncAlignedShape2NumUnpackB32 {
266        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
267            stream.expect_string("tcgen05")?;
268            stream.expect_string(".st")?;
269            let st = ();
270            stream.expect_complete()?;
271            stream.expect_string(".sync")?;
272            let sync = ();
273            stream.expect_complete()?;
274            stream.expect_string(".aligned")?;
275            let aligned = ();
276            stream.expect_complete()?;
277            let shape2 = Shape2::parse(stream)?;
278            stream.expect_complete()?;
279            let num = Num::parse(stream)?;
280            stream.expect_complete()?;
281            let saved_pos = stream.position();
282            let unpack = match Unpack::parse(stream) {
283                Ok(val) => Some(val),
284                Err(_) => {
285                    stream.set_position(saved_pos);
286                    None
287                }
288            };
289            stream.expect_complete()?;
290            stream.expect_string(".b32")?;
291            let b32 = ();
292            stream.expect_complete()?;
293            let taddr = AddressOperand::parse(stream)?;
294            stream.expect_complete()?;
295            stream.expect(&PtxToken::Comma)?;
296            let immhalfsplitoff = GeneralOperand::parse(stream)?;
297            stream.expect_complete()?;
298            stream.expect(&PtxToken::Comma)?;
299            let r = GeneralOperand::parse(stream)?;
300            stream.expect_complete()?;
301            stream.expect_complete()?;
302            stream.expect(&PtxToken::Semicolon)?;
303            Ok(Tcgen05StSyncAlignedShape2NumUnpackB32 {
304                st,
305                sync,
306                aligned,
307                shape2,
308                num,
309                unpack,
310                b32,
311                taddr,
312                immhalfsplitoff,
313                r,
314            })
315        }
316    }
317}
ptx_parser/parser/instruction/tcgen05_st.rs

ptx_parser/parser/instruction/
tcgen05_st.rs
