ptx_parser/parser/instruction/
stmatrix.rs

1//! Original PTX specification:
2//!
3//! stmatrix.sync.aligned.shape.num{.trans}{.ss}.type [p], r;
4//! .shape  = {.m8n8, .m16n8};
5//! .num    = {.x1, .x2, .x4};
6//! .ss     = {.shared, .shared::cta};
7//! .type   = {.b16, .b8};
8
9#![allow(unused)]
10
11use crate::lexer::PtxToken;
12use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
13use crate::r#type::common::*;
14
15pub mod section_0 {
16    use super::*;
17    use crate::r#type::instruction::stmatrix::section_0::*;
18
19    // ============================================================================
20    // Generated enum parsers
21    // ============================================================================
22
23    impl PtxParser for Num {
24        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
25            // Try X1
26            {
27                let saved_pos = stream.position();
28                if stream.expect_string(".x1").is_ok() {
29                    return Ok(Num::X1);
30                }
31                stream.set_position(saved_pos);
32            }
33            let saved_pos = stream.position();
34            // Try X2
35            {
36                let saved_pos = stream.position();
37                if stream.expect_string(".x2").is_ok() {
38                    return Ok(Num::X2);
39                }
40                stream.set_position(saved_pos);
41            }
42            stream.set_position(saved_pos);
43            let saved_pos = stream.position();
44            // Try X4
45            {
46                let saved_pos = stream.position();
47                if stream.expect_string(".x4").is_ok() {
48                    return Ok(Num::X4);
49                }
50                stream.set_position(saved_pos);
51            }
52            stream.set_position(saved_pos);
53            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
54            let expected = &[".x1", ".x2", ".x4"];
55            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
56            Err(crate::parser::unexpected_value(span, expected, found))
57        }
58    }
59
60    impl PtxParser for Shape {
61        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
62            // Try M16n8
63            {
64                let saved_pos = stream.position();
65                if stream.expect_string(".m16n8").is_ok() {
66                    return Ok(Shape::M16n8);
67                }
68                stream.set_position(saved_pos);
69            }
70            let saved_pos = stream.position();
71            // Try M8n8
72            {
73                let saved_pos = stream.position();
74                if stream.expect_string(".m8n8").is_ok() {
75                    return Ok(Shape::M8n8);
76                }
77                stream.set_position(saved_pos);
78            }
79            stream.set_position(saved_pos);
80            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
81            let expected = &[".m16n8", ".m8n8"];
82            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
83            Err(crate::parser::unexpected_value(span, expected, found))
84        }
85    }
86
87    impl PtxParser for Ss {
88        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
89            // Try SharedCta
90            {
91                let saved_pos = stream.position();
92                if stream.expect_string(".shared::cta").is_ok() {
93                    return Ok(Ss::SharedCta);
94                }
95                stream.set_position(saved_pos);
96            }
97            let saved_pos = stream.position();
98            // Try Shared
99            {
100                let saved_pos = stream.position();
101                if stream.expect_string(".shared").is_ok() {
102                    return Ok(Ss::Shared);
103                }
104                stream.set_position(saved_pos);
105            }
106            stream.set_position(saved_pos);
107            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
108            let expected = &[".shared::cta", ".shared"];
109            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
110            Err(crate::parser::unexpected_value(span, expected, found))
111        }
112    }
113
114    impl PtxParser for Type {
115        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
116            // Try B16
117            {
118                let saved_pos = stream.position();
119                if stream.expect_string(".b16").is_ok() {
120                    return Ok(Type::B16);
121                }
122                stream.set_position(saved_pos);
123            }
124            let saved_pos = stream.position();
125            // Try B8
126            {
127                let saved_pos = stream.position();
128                if stream.expect_string(".b8").is_ok() {
129                    return Ok(Type::B8);
130                }
131                stream.set_position(saved_pos);
132            }
133            stream.set_position(saved_pos);
134            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
135            let expected = &[".b16", ".b8"];
136            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
137            Err(crate::parser::unexpected_value(span, expected, found))
138        }
139    }
140
141    impl PtxParser for StmatrixSyncAlignedShapeNumTransSsType {
142        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
143            stream.expect_string("stmatrix")?;
144            stream.expect_string(".sync")?;
145            let sync = ();
146            stream.expect_complete()?;
147            stream.expect_string(".aligned")?;
148            let aligned = ();
149            stream.expect_complete()?;
150            let shape = Shape::parse(stream)?;
151            stream.expect_complete()?;
152            let num = Num::parse(stream)?;
153            stream.expect_complete()?;
154            let saved_pos = stream.position();
155            let trans = stream.expect_string(".trans").is_ok();
156            if !trans {
157                stream.set_position(saved_pos);
158            }
159            stream.expect_complete()?;
160            let saved_pos = stream.position();
161            let ss = match Ss::parse(stream) {
162                Ok(val) => Some(val),
163                Err(_) => {
164                    stream.set_position(saved_pos);
165                    None
166                }
167            };
168            stream.expect_complete()?;
169            let type_ = Type::parse(stream)?;
170            stream.expect_complete()?;
171            let p = AddressOperand::parse(stream)?;
172            stream.expect_complete()?;
173            stream.expect(&PtxToken::Comma)?;
174            let r = GeneralOperand::parse(stream)?;
175            stream.expect_complete()?;
176            stream.expect_complete()?;
177            stream.expect(&PtxToken::Semicolon)?;
178            Ok(StmatrixSyncAlignedShapeNumTransSsType {
179                sync,
180                aligned,
181                shape,
182                num,
183                trans,
184                ss,
185                type_,
186                p,
187                r,
188            })
189        }
190    }
191
192
193}
194
ptx_parser/parser/instruction/stmatrix.rs

ptx_parser/parser/instruction/
stmatrix.rs
