ptx_parser/parser/instruction/
tcgen05_cp.rs

1//! Original PTX specification:
2//!
3//! tcgen05.cp.cta_group.shape{.multicast}{.dst_src_fmt} [taddr], s-desc;
4//! .cta_group = { .cta_group::1, .cta_group::2 };
5//! .dst_src_fmt   = { .b8x16.b6x16_p32, .b8x16.b4x16_p64 };
6//! .shape     = { .128x256b, .4x256b, .128x128b, .64x128b**, .32x128b*** };
7//! .multicast = { .warpx2::02_13** , .warpx2::01_23**, .warpx4*** };
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::tcgen05_cp::section_0::*;
18
19    // ============================================================================
20    // Generated enum parsers
21    // ============================================================================
22
23    impl PtxParser for CtaGroup {
24        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
25            // Try CtaGroup1
26            {
27                let saved_pos = stream.position();
28                if stream.expect_string(".cta_group::1").is_ok() {
29                    return Ok(CtaGroup::CtaGroup1);
30                }
31                stream.set_position(saved_pos);
32            }
33            let saved_pos = stream.position();
34            // Try CtaGroup2
35            {
36                let saved_pos = stream.position();
37                if stream.expect_string(".cta_group::2").is_ok() {
38                    return Ok(CtaGroup::CtaGroup2);
39                }
40                stream.set_position(saved_pos);
41            }
42            stream.set_position(saved_pos);
43            let span = stream
44                .peek()
45                .map(|(_, s)| s.clone())
46                .unwrap_or(Span { start: 0, end: 0 });
47            let expected = &[".cta_group::1", ".cta_group::2"];
48            let found = stream
49                .peek()
50                .map(|(t, _)| format!("{:?}", t))
51                .unwrap_or_else(|_| "<end of input>".to_string());
52            Err(crate::parser::unexpected_value(span, expected, found))
53        }
54    }
55
56    impl PtxParser for DstSrcFmt {
57        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
58            // Try B8x16B6x16P32
59            {
60                let saved_pos = stream.position();
61                if stream.expect_string(".b8x16.b6x16_p32").is_ok() {
62                    return Ok(DstSrcFmt::B8x16B6x16P32);
63                }
64                stream.set_position(saved_pos);
65            }
66            let saved_pos = stream.position();
67            // Try B8x16B4x16P64
68            {
69                let saved_pos = stream.position();
70                if stream.expect_string(".b8x16.b4x16_p64").is_ok() {
71                    return Ok(DstSrcFmt::B8x16B4x16P64);
72                }
73                stream.set_position(saved_pos);
74            }
75            stream.set_position(saved_pos);
76            let span = stream
77                .peek()
78                .map(|(_, s)| s.clone())
79                .unwrap_or(Span { start: 0, end: 0 });
80            let expected = &[".b8x16.b6x16_p32", ".b8x16.b4x16_p64"];
81            let found = stream
82                .peek()
83                .map(|(t, _)| format!("{:?}", t))
84                .unwrap_or_else(|_| "<end of input>".to_string());
85            Err(crate::parser::unexpected_value(span, expected, found))
86        }
87    }
88
89    impl PtxParser for Multicast {
90        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
91            // Try Warpx20213
92            {
93                let saved_pos = stream.position();
94                if stream.expect_string(".warpx2::02_13**").is_ok() {
95                    return Ok(Multicast::Warpx20213);
96                }
97                stream.set_position(saved_pos);
98            }
99            let saved_pos = stream.position();
100            // Try Warpx20123
101            {
102                let saved_pos = stream.position();
103                if stream.expect_string(".warpx2::01_23**").is_ok() {
104                    return Ok(Multicast::Warpx20123);
105                }
106                stream.set_position(saved_pos);
107            }
108            stream.set_position(saved_pos);
109            let saved_pos = stream.position();
110            // Try Warpx4
111            {
112                let saved_pos = stream.position();
113                if stream.expect_string(".warpx4***").is_ok() {
114                    return Ok(Multicast::Warpx4);
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 = &[".warpx2::02_13**", ".warpx2::01_23**", ".warpx4***"];
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 Shape {
133        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
134            // Try _32x128b
135            {
136                let saved_pos = stream.position();
137                if stream.expect_string(".32x128b***").is_ok() {
138                    return Ok(Shape::_32x128b);
139                }
140                stream.set_position(saved_pos);
141            }
142            let saved_pos = stream.position();
143            // Try _64x128b
144            {
145                let saved_pos = stream.position();
146                if stream.expect_string(".64x128b**").is_ok() {
147                    return Ok(Shape::_64x128b);
148                }
149                stream.set_position(saved_pos);
150            }
151            stream.set_position(saved_pos);
152            let saved_pos = stream.position();
153            // Try _128x256b
154            {
155                let saved_pos = stream.position();
156                if stream.expect_string(".128x256b").is_ok() {
157                    return Ok(Shape::_128x256b);
158                }
159                stream.set_position(saved_pos);
160            }
161            stream.set_position(saved_pos);
162            let saved_pos = stream.position();
163            // Try _128x128b
164            {
165                let saved_pos = stream.position();
166                if stream.expect_string(".128x128b").is_ok() {
167                    return Ok(Shape::_128x128b);
168                }
169                stream.set_position(saved_pos);
170            }
171            stream.set_position(saved_pos);
172            let saved_pos = stream.position();
173            // Try _4x256b
174            {
175                let saved_pos = stream.position();
176                if stream.expect_string(".4x256b").is_ok() {
177                    return Ok(Shape::_4x256b);
178                }
179                stream.set_position(saved_pos);
180            }
181            stream.set_position(saved_pos);
182            let span = stream
183                .peek()
184                .map(|(_, s)| s.clone())
185                .unwrap_or(Span { start: 0, end: 0 });
186            let expected = &[
187                ".32x128b***",
188                ".64x128b**",
189                ".128x256b",
190                ".128x128b",
191                ".4x256b",
192            ];
193            let found = stream
194                .peek()
195                .map(|(t, _)| format!("{:?}", t))
196                .unwrap_or_else(|_| "<end of input>".to_string());
197            Err(crate::parser::unexpected_value(span, expected, found))
198        }
199    }
200
201    impl PtxParser for Tcgen05CpCtaGroupShapeMulticastDstSrcFmt {
202        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
203            stream.expect_string("tcgen05")?;
204            stream.expect_string(".cp")?;
205            let cp = ();
206            stream.expect_complete()?;
207            let cta_group = CtaGroup::parse(stream)?;
208            stream.expect_complete()?;
209            let shape = Shape::parse(stream)?;
210            stream.expect_complete()?;
211            let saved_pos = stream.position();
212            let multicast = match Multicast::parse(stream) {
213                Ok(val) => Some(val),
214                Err(_) => {
215                    stream.set_position(saved_pos);
216                    None
217                }
218            };
219            stream.expect_complete()?;
220            let saved_pos = stream.position();
221            let dst_src_fmt = match DstSrcFmt::parse(stream) {
222                Ok(val) => Some(val),
223                Err(_) => {
224                    stream.set_position(saved_pos);
225                    None
226                }
227            };
228            stream.expect_complete()?;
229            let taddr = AddressOperand::parse(stream)?;
230            stream.expect_complete()?;
231            stream.expect(&PtxToken::Comma)?;
232            let s_desc = GeneralOperand::parse(stream)?;
233            stream.expect_complete()?;
234            stream.expect_complete()?;
235            stream.expect(&PtxToken::Semicolon)?;
236            Ok(Tcgen05CpCtaGroupShapeMulticastDstSrcFmt {
237                cp,
238                cta_group,
239                shape,
240                multicast,
241                dst_src_fmt,
242                taddr,
243                s_desc,
244            })
245        }
246    }
247}
ptx_parser/parser/instruction/tcgen05_cp.rs

ptx_parser/parser/instruction/
tcgen05_cp.rs
