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

ptx_parser/parser/instruction/
ldmatrix.rs
