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.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
44            let expected = &[".cta_group::1", ".cta_group::2"];
45            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
46            Err(crate::parser::unexpected_value(span, expected, found))
47        }
48    }
49
50    impl PtxParser for DstSrcFmt {
51        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
52            // Try B8x16B6x16P32
53            {
54                let saved_pos = stream.position();
55                if stream.expect_string(".b8x16.b6x16_p32").is_ok() {
56                    return Ok(DstSrcFmt::B8x16B6x16P32);
57                }
58                stream.set_position(saved_pos);
59            }
60            let saved_pos = stream.position();
61            // Try B8x16B4x16P64
62            {
63                let saved_pos = stream.position();
64                if stream.expect_string(".b8x16.b4x16_p64").is_ok() {
65                    return Ok(DstSrcFmt::B8x16B4x16P64);
66                }
67                stream.set_position(saved_pos);
68            }
69            stream.set_position(saved_pos);
70            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
71            let expected = &[".b8x16.b6x16_p32", ".b8x16.b4x16_p64"];
72            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
73            Err(crate::parser::unexpected_value(span, expected, found))
74        }
75    }
76
77    impl PtxParser for Multicast {
78        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
79            // Try Warpx20213
80            {
81                let saved_pos = stream.position();
82                if stream.expect_string(".warpx2::02_13**").is_ok() {
83                    return Ok(Multicast::Warpx20213);
84                }
85                stream.set_position(saved_pos);
86            }
87            let saved_pos = stream.position();
88            // Try Warpx20123
89            {
90                let saved_pos = stream.position();
91                if stream.expect_string(".warpx2::01_23**").is_ok() {
92                    return Ok(Multicast::Warpx20123);
93                }
94                stream.set_position(saved_pos);
95            }
96            stream.set_position(saved_pos);
97            let saved_pos = stream.position();
98            // Try Warpx4
99            {
100                let saved_pos = stream.position();
101                if stream.expect_string(".warpx4***").is_ok() {
102                    return Ok(Multicast::Warpx4);
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 = &[".warpx2::02_13**", ".warpx2::01_23**", ".warpx4***"];
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 Shape {
115        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
116            // Try _32x128b
117            {
118                let saved_pos = stream.position();
119                if stream.expect_string(".32x128b***").is_ok() {
120                    return Ok(Shape::_32x128b);
121                }
122                stream.set_position(saved_pos);
123            }
124            let saved_pos = stream.position();
125            // Try _64x128b
126            {
127                let saved_pos = stream.position();
128                if stream.expect_string(".64x128b**").is_ok() {
129                    return Ok(Shape::_64x128b);
130                }
131                stream.set_position(saved_pos);
132            }
133            stream.set_position(saved_pos);
134            let saved_pos = stream.position();
135            // Try _128x256b
136            {
137                let saved_pos = stream.position();
138                if stream.expect_string(".128x256b").is_ok() {
139                    return Ok(Shape::_128x256b);
140                }
141                stream.set_position(saved_pos);
142            }
143            stream.set_position(saved_pos);
144            let saved_pos = stream.position();
145            // Try _128x128b
146            {
147                let saved_pos = stream.position();
148                if stream.expect_string(".128x128b").is_ok() {
149                    return Ok(Shape::_128x128b);
150                }
151                stream.set_position(saved_pos);
152            }
153            stream.set_position(saved_pos);
154            let saved_pos = stream.position();
155            // Try _4x256b
156            {
157                let saved_pos = stream.position();
158                if stream.expect_string(".4x256b").is_ok() {
159                    return Ok(Shape::_4x256b);
160                }
161                stream.set_position(saved_pos);
162            }
163            stream.set_position(saved_pos);
164            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
165            let expected = &[".32x128b***", ".64x128b**", ".128x256b", ".128x128b", ".4x256b"];
166            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
167            Err(crate::parser::unexpected_value(span, expected, found))
168        }
169    }
170
171    impl PtxParser for Tcgen05CpCtaGroupShapeMulticastDstSrcFmt {
172        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
173            stream.expect_string("tcgen05")?;
174            stream.expect_string(".cp")?;
175            let cp = ();
176            stream.expect_complete()?;
177            let cta_group = CtaGroup::parse(stream)?;
178            stream.expect_complete()?;
179            let shape = Shape::parse(stream)?;
180            stream.expect_complete()?;
181            let saved_pos = stream.position();
182            let multicast = match Multicast::parse(stream) {
183                Ok(val) => Some(val),
184                Err(_) => {
185                    stream.set_position(saved_pos);
186                    None
187                }
188            };
189            stream.expect_complete()?;
190            let saved_pos = stream.position();
191            let dst_src_fmt = match DstSrcFmt::parse(stream) {
192                Ok(val) => Some(val),
193                Err(_) => {
194                    stream.set_position(saved_pos);
195                    None
196                }
197            };
198            stream.expect_complete()?;
199            let taddr = AddressOperand::parse(stream)?;
200            stream.expect_complete()?;
201            stream.expect(&PtxToken::Comma)?;
202            let s_desc = GeneralOperand::parse(stream)?;
203            stream.expect_complete()?;
204            stream.expect_complete()?;
205            stream.expect(&PtxToken::Semicolon)?;
206            Ok(Tcgen05CpCtaGroupShapeMulticastDstSrcFmt {
207                cp,
208                cta_group,
209                shape,
210                multicast,
211                dst_src_fmt,
212                taddr,
213                s_desc,
214            })
215        }
216    }
217
218
219}
220
ptx_parser/parser/instruction/tcgen05_cp.rs

ptx_parser/parser/instruction/
tcgen05_cp.rs
