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
54                .peek()
55                .map(|(_, s)| s.clone())
56                .unwrap_or(Span { start: 0, end: 0 });
57            let expected = &[".x1", ".x2", ".x4"];
58            let found = stream
59                .peek()
60                .map(|(t, _)| format!("{:?}", t))
61                .unwrap_or_else(|_| "<end of input>".to_string());
62            Err(crate::parser::unexpected_value(span, expected, found))
63        }
64    }
65
66    impl PtxParser for Shape {
67        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
68            // Try M16n8
69            {
70                let saved_pos = stream.position();
71                if stream.expect_string(".m16n8").is_ok() {
72                    return Ok(Shape::M16n8);
73                }
74                stream.set_position(saved_pos);
75            }
76            let saved_pos = stream.position();
77            // Try M8n8
78            {
79                let saved_pos = stream.position();
80                if stream.expect_string(".m8n8").is_ok() {
81                    return Ok(Shape::M8n8);
82                }
83                stream.set_position(saved_pos);
84            }
85            stream.set_position(saved_pos);
86            let span = stream
87                .peek()
88                .map(|(_, s)| s.clone())
89                .unwrap_or(Span { start: 0, end: 0 });
90            let expected = &[".m16n8", ".m8n8"];
91            let found = stream
92                .peek()
93                .map(|(t, _)| format!("{:?}", t))
94                .unwrap_or_else(|_| "<end of input>".to_string());
95            Err(crate::parser::unexpected_value(span, expected, found))
96        }
97    }
98
99    impl PtxParser for Ss {
100        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
101            // Try SharedCta
102            {
103                let saved_pos = stream.position();
104                if stream.expect_string(".shared::cta").is_ok() {
105                    return Ok(Ss::SharedCta);
106                }
107                stream.set_position(saved_pos);
108            }
109            let saved_pos = stream.position();
110            // Try Shared
111            {
112                let saved_pos = stream.position();
113                if stream.expect_string(".shared").is_ok() {
114                    return Ok(Ss::Shared);
115                }
116                stream.set_position(saved_pos);
117            }
118            stream.set_position(saved_pos);
119            let span = stream
120                .peek()
121                .map(|(_, s)| s.clone())
122                .unwrap_or(Span { start: 0, end: 0 });
123            let expected = &[".shared::cta", ".shared"];
124            let found = stream
125                .peek()
126                .map(|(t, _)| format!("{:?}", t))
127                .unwrap_or_else(|_| "<end of input>".to_string());
128            Err(crate::parser::unexpected_value(span, expected, found))
129        }
130    }
131
132    impl PtxParser for Type {
133        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
134            // Try B16
135            {
136                let saved_pos = stream.position();
137                if stream.expect_string(".b16").is_ok() {
138                    return Ok(Type::B16);
139                }
140                stream.set_position(saved_pos);
141            }
142            let saved_pos = stream.position();
143            // Try B8
144            {
145                let saved_pos = stream.position();
146                if stream.expect_string(".b8").is_ok() {
147                    return Ok(Type::B8);
148                }
149                stream.set_position(saved_pos);
150            }
151            stream.set_position(saved_pos);
152            let span = stream
153                .peek()
154                .map(|(_, s)| s.clone())
155                .unwrap_or(Span { start: 0, end: 0 });
156            let expected = &[".b16", ".b8"];
157            let found = stream
158                .peek()
159                .map(|(t, _)| format!("{:?}", t))
160                .unwrap_or_else(|_| "<end of input>".to_string());
161            Err(crate::parser::unexpected_value(span, expected, found))
162        }
163    }
164
165    impl PtxParser for StmatrixSyncAlignedShapeNumTransSsType {
166        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
167            stream.expect_string("stmatrix")?;
168            stream.expect_string(".sync")?;
169            let sync = ();
170            stream.expect_complete()?;
171            stream.expect_string(".aligned")?;
172            let aligned = ();
173            stream.expect_complete()?;
174            let shape = Shape::parse(stream)?;
175            stream.expect_complete()?;
176            let num = Num::parse(stream)?;
177            stream.expect_complete()?;
178            let saved_pos = stream.position();
179            let trans = stream.expect_string(".trans").is_ok();
180            if !trans {
181                stream.set_position(saved_pos);
182            }
183            stream.expect_complete()?;
184            let saved_pos = stream.position();
185            let ss = match Ss::parse(stream) {
186                Ok(val) => Some(val),
187                Err(_) => {
188                    stream.set_position(saved_pos);
189                    None
190                }
191            };
192            stream.expect_complete()?;
193            let type_ = Type::parse(stream)?;
194            stream.expect_complete()?;
195            let p = AddressOperand::parse(stream)?;
196            stream.expect_complete()?;
197            stream.expect(&PtxToken::Comma)?;
198            let r = GeneralOperand::parse(stream)?;
199            stream.expect_complete()?;
200            stream.expect_complete()?;
201            stream.expect(&PtxToken::Semicolon)?;
202            Ok(StmatrixSyncAlignedShapeNumTransSsType {
203                sync,
204                aligned,
205                shape,
206                num,
207                trans,
208                ss,
209                type_,
210                p,
211                r,
212            })
213        }
214    }
215}
ptx_parser/parser/instruction/stmatrix.rs

ptx_parser/parser/instruction/
stmatrix.rs
