ptx_parser/parser/instruction/
ldmatrix.rs

1//! Original PTX specification:
2//!
3//! ldmatrix.sync.aligned.shape.num{.trans}{.ss}.type r, [p];
4//! ldmatrix.sync.aligned.m8n16.num{.ss}.dst_fmt.src_fmt        r, [p];
5//! ldmatrix.sync.aligned.m16n16.num.trans{.ss}.dst_fmt.src_fmt r, [p];
6//! .shape   = {.m8n8, .m16n16};
7//! .num     = {.x1, .x2, .x4};
8//! .ss      = {.shared, .shared::cta};
9//! .type    = {.b16, .b8};
10//! .dst_fmt = { .b8x16 };
11//! .src_fmt = { .b6x16_p32, .b4x16_p64 };
12
13#![allow(unused)]
14
15use crate::lexer::PtxToken;
16use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
17use crate::r#type::common::*;
18
19pub mod section_0 {
20    use super::*;
21    use crate::r#type::instruction::ldmatrix::section_0::*;
22
23    // ============================================================================
24    // Generated enum parsers
25    // ============================================================================
26
27    impl PtxParser for DstFmt {
28        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
29            // Try B8x16
30            {
31                let saved_pos = stream.position();
32                if stream.expect_string(".b8x16").is_ok() {
33                    return Ok(DstFmt::B8x16);
34                }
35                stream.set_position(saved_pos);
36            }
37            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
38            let expected = &[".b8x16"];
39            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
40            Err(crate::parser::unexpected_value(span, expected, found))
41        }
42    }
43
44    impl PtxParser for Num {
45        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
46            // Try X1
47            {
48                let saved_pos = stream.position();
49                if stream.expect_string(".x1").is_ok() {
50                    return Ok(Num::X1);
51                }
52                stream.set_position(saved_pos);
53            }
54            let saved_pos = stream.position();
55            // Try X2
56            {
57                let saved_pos = stream.position();
58                if stream.expect_string(".x2").is_ok() {
59                    return Ok(Num::X2);
60                }
61                stream.set_position(saved_pos);
62            }
63            stream.set_position(saved_pos);
64            let saved_pos = stream.position();
65            // Try X4
66            {
67                let saved_pos = stream.position();
68                if stream.expect_string(".x4").is_ok() {
69                    return Ok(Num::X4);
70                }
71                stream.set_position(saved_pos);
72            }
73            stream.set_position(saved_pos);
74            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
75            let expected = &[".x1", ".x2", ".x4"];
76            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
77            Err(crate::parser::unexpected_value(span, expected, found))
78        }
79    }
80
81    impl PtxParser for Shape {
82        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
83            // Try M16n16
84            {
85                let saved_pos = stream.position();
86                if stream.expect_string(".m16n16").is_ok() {
87                    return Ok(Shape::M16n16);
88                }
89                stream.set_position(saved_pos);
90            }
91            let saved_pos = stream.position();
92            // Try M8n8
93            {
94                let saved_pos = stream.position();
95                if stream.expect_string(".m8n8").is_ok() {
96                    return Ok(Shape::M8n8);
97                }
98                stream.set_position(saved_pos);
99            }
100            stream.set_position(saved_pos);
101            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
102            let expected = &[".m16n16", ".m8n8"];
103            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
104            Err(crate::parser::unexpected_value(span, expected, found))
105        }
106    }
107
108    impl PtxParser for SrcFmt {
109        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
110            // Try B6x16P32
111            {
112                let saved_pos = stream.position();
113                if stream.expect_string(".b6x16_p32").is_ok() {
114                    return Ok(SrcFmt::B6x16P32);
115                }
116                stream.set_position(saved_pos);
117            }
118            let saved_pos = stream.position();
119            // Try B4x16P64
120            {
121                let saved_pos = stream.position();
122                if stream.expect_string(".b4x16_p64").is_ok() {
123                    return Ok(SrcFmt::B4x16P64);
124                }
125                stream.set_position(saved_pos);
126            }
127            stream.set_position(saved_pos);
128            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
129            let expected = &[".b6x16_p32", ".b4x16_p64"];
130            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
131            Err(crate::parser::unexpected_value(span, expected, found))
132        }
133    }
134
135    impl PtxParser for Ss {
136        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
137            // Try SharedCta
138            {
139                let saved_pos = stream.position();
140                if stream.expect_string(".shared::cta").is_ok() {
141                    return Ok(Ss::SharedCta);
142                }
143                stream.set_position(saved_pos);
144            }
145            let saved_pos = stream.position();
146            // Try Shared
147            {
148                let saved_pos = stream.position();
149                if stream.expect_string(".shared").is_ok() {
150                    return Ok(Ss::Shared);
151                }
152                stream.set_position(saved_pos);
153            }
154            stream.set_position(saved_pos);
155            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
156            let expected = &[".shared::cta", ".shared"];
157            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
158            Err(crate::parser::unexpected_value(span, expected, found))
159        }
160    }
161
162    impl PtxParser for Type {
163        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
164            // Try B16
165            {
166                let saved_pos = stream.position();
167                if stream.expect_string(".b16").is_ok() {
168                    return Ok(Type::B16);
169                }
170                stream.set_position(saved_pos);
171            }
172            let saved_pos = stream.position();
173            // Try B8
174            {
175                let saved_pos = stream.position();
176                if stream.expect_string(".b8").is_ok() {
177                    return Ok(Type::B8);
178                }
179                stream.set_position(saved_pos);
180            }
181            stream.set_position(saved_pos);
182            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
183            let expected = &[".b16", ".b8"];
184            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
185            Err(crate::parser::unexpected_value(span, expected, found))
186        }
187    }
188
189    impl PtxParser for LdmatrixSyncAlignedShapeNumTransSsType {
190        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
191            stream.expect_string("ldmatrix")?;
192            stream.expect_string(".sync")?;
193            let sync = ();
194            stream.expect_complete()?;
195            stream.expect_string(".aligned")?;
196            let aligned = ();
197            stream.expect_complete()?;
198            let shape = Shape::parse(stream)?;
199            stream.expect_complete()?;
200            let num = Num::parse(stream)?;
201            stream.expect_complete()?;
202            let saved_pos = stream.position();
203            let trans = stream.expect_string(".trans").is_ok();
204            if !trans {
205                stream.set_position(saved_pos);
206            }
207            stream.expect_complete()?;
208            let saved_pos = stream.position();
209            let ss = match Ss::parse(stream) {
210                Ok(val) => Some(val),
211                Err(_) => {
212                    stream.set_position(saved_pos);
213                    None
214                }
215            };
216            stream.expect_complete()?;
217            let type_ = Type::parse(stream)?;
218            stream.expect_complete()?;
219            let r = GeneralOperand::parse(stream)?;
220            stream.expect_complete()?;
221            stream.expect(&PtxToken::Comma)?;
222            let p = AddressOperand::parse(stream)?;
223            stream.expect_complete()?;
224            stream.expect_complete()?;
225            stream.expect(&PtxToken::Semicolon)?;
226            Ok(LdmatrixSyncAlignedShapeNumTransSsType {
227                sync,
228                aligned,
229                shape,
230                num,
231                trans,
232                ss,
233                type_,
234                r,
235                p,
236            })
237        }
238    }
239
240
241    impl PtxParser for LdmatrixSyncAlignedM8n16NumSsDstFmtSrcFmt {
242        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
243            stream.expect_string("ldmatrix")?;
244            stream.expect_string(".sync")?;
245            let sync = ();
246            stream.expect_complete()?;
247            stream.expect_string(".aligned")?;
248            let aligned = ();
249            stream.expect_complete()?;
250            stream.expect_string(".m8n16")?;
251            let m8n16 = ();
252            stream.expect_complete()?;
253            let num = Num::parse(stream)?;
254            stream.expect_complete()?;
255            let saved_pos = stream.position();
256            let ss = match Ss::parse(stream) {
257                Ok(val) => Some(val),
258                Err(_) => {
259                    stream.set_position(saved_pos);
260                    None
261                }
262            };
263            stream.expect_complete()?;
264            let dst_fmt = DstFmt::parse(stream)?;
265            stream.expect_complete()?;
266            let src_fmt = SrcFmt::parse(stream)?;
267            stream.expect_complete()?;
268            let r = GeneralOperand::parse(stream)?;
269            stream.expect_complete()?;
270            stream.expect(&PtxToken::Comma)?;
271            let p = AddressOperand::parse(stream)?;
272            stream.expect_complete()?;
273            stream.expect_complete()?;
274            stream.expect(&PtxToken::Semicolon)?;
275            Ok(LdmatrixSyncAlignedM8n16NumSsDstFmtSrcFmt {
276                sync,
277                aligned,
278                m8n16,
279                num,
280                ss,
281                dst_fmt,
282                src_fmt,
283                r,
284                p,
285            })
286        }
287    }
288
289
290    impl PtxParser for LdmatrixSyncAlignedM16n16NumTransSsDstFmtSrcFmt {
291        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
292            stream.expect_string("ldmatrix")?;
293            stream.expect_string(".sync")?;
294            let sync = ();
295            stream.expect_complete()?;
296            stream.expect_string(".aligned")?;
297            let aligned = ();
298            stream.expect_complete()?;
299            stream.expect_string(".m16n16")?;
300            let m16n16 = ();
301            stream.expect_complete()?;
302            let num = Num::parse(stream)?;
303            stream.expect_complete()?;
304            stream.expect_string(".trans")?;
305            let trans = ();
306            stream.expect_complete()?;
307            let saved_pos = stream.position();
308            let ss = match Ss::parse(stream) {
309                Ok(val) => Some(val),
310                Err(_) => {
311                    stream.set_position(saved_pos);
312                    None
313                }
314            };
315            stream.expect_complete()?;
316            let dst_fmt = DstFmt::parse(stream)?;
317            stream.expect_complete()?;
318            let src_fmt = SrcFmt::parse(stream)?;
319            stream.expect_complete()?;
320            let r = GeneralOperand::parse(stream)?;
321            stream.expect_complete()?;
322            stream.expect(&PtxToken::Comma)?;
323            let p = AddressOperand::parse(stream)?;
324            stream.expect_complete()?;
325            stream.expect_complete()?;
326            stream.expect(&PtxToken::Semicolon)?;
327            Ok(LdmatrixSyncAlignedM16n16NumTransSsDstFmtSrcFmt {
328                sync,
329                aligned,
330                m16n16,
331                num,
332                trans,
333                ss,
334                dst_fmt,
335                src_fmt,
336                r,
337                p,
338            })
339        }
340    }
341
342
343}
344
ptx_parser/parser/instruction/ldmatrix.rs

ptx_parser/parser/instruction/
ldmatrix.rs
