ptx_parser/parser/instruction/
tcgen05_mma_ws.rs

1//! Original PTX specification:
2//!
3//! // 1. Floating-point type without block scaling:
4//! tcgen05.mma.ws.cta_group::1.kind{.collector_usage}    [d-tmem],  a-desc,  b-desc,  idesc,
5//! enable-input-d {, zero-column-mask-desc };
6//! tcgen05.mma.ws.cta_group::1.kind{.collector_usage}    [d-tmem], [a-tmem], b-desc, idesc,
7//! enable-input-d {, zero-column-mask-desc };
8//! .kind = { .kind::f16, .kind::tf32, .kind::f8f6f4 };
9//! ----------------------------------------------------------------------------------
10//! // 2. Integer type:
11//! tcgen05.mma.ws.cta_group::1.kind::i8{.collector_usage} [d-tmem],  a-desc,  b-desc, idesc,
12//! enable-input-d {, zero-column-mask-desc};
13//! tcgen05.mma.ws.cta_group::1.kind::i8{.collector_usage} [d-tmem], [a-tmem], b-desc, idesc,
14//! enable-input-d {, zero-column-mask-desc};
15//! .collector_usage = { .collector::buffer::op };
16//! ::buffer = { ::b0, ::b1, ::b2, ::b3 };
17//! ::op   = { ::fill, ::use, ::lastuse, ::discard};
18
19#![allow(unused)]
20
21use crate::lexer::PtxToken;
22use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
23use crate::r#type::common::*;
24
25pub mod section_0 {
26    use super::*;
27    use crate::r#type::instruction::tcgen05_mma_ws::section_0::*;
28
29    // ============================================================================
30    // Generated enum parsers
31    // ============================================================================
32
33    impl PtxParser for Kind {
34        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
35            // Try KindF8f6f4
36            {
37                let saved_pos = stream.position();
38                if stream.expect_string(".kind::f8f6f4").is_ok() {
39                    return Ok(Kind::KindF8f6f4);
40                }
41                stream.set_position(saved_pos);
42            }
43            let saved_pos = stream.position();
44            // Try KindTf32
45            {
46                let saved_pos = stream.position();
47                if stream.expect_string(".kind::tf32").is_ok() {
48                    return Ok(Kind::KindTf32);
49                }
50                stream.set_position(saved_pos);
51            }
52            stream.set_position(saved_pos);
53            let saved_pos = stream.position();
54            // Try KindF16
55            {
56                let saved_pos = stream.position();
57                if stream.expect_string(".kind::f16").is_ok() {
58                    return Ok(Kind::KindF16);
59                }
60                stream.set_position(saved_pos);
61            }
62            stream.set_position(saved_pos);
63            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
64            let expected = &[".kind::f8f6f4", ".kind::tf32", ".kind::f16"];
65            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
66            Err(crate::parser::unexpected_value(span, expected, found))
67        }
68    }
69
70    impl PtxParser for Tcgen05MmaWsCtaGroup1KindCollectorUsage {
71        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
72            stream.expect_string("tcgen05")?;
73            stream.expect_string(".mma")?;
74            let mma = ();
75            stream.expect_complete()?;
76            stream.expect_string(".ws")?;
77            let ws = ();
78            stream.expect_complete()?;
79            stream.expect_string(".cta_group::1")?;
80            let cta_group_1 = ();
81            stream.expect_complete()?;
82            let kind = Kind::parse(stream)?;
83            stream.expect_complete()?;
84            let saved_pos = stream.position();
85            let collector_usage = stream.expect_string(".collector_usage").is_ok();
86            if !collector_usage {
87                stream.set_position(saved_pos);
88            }
89            stream.expect_complete()?;
90            let d_tmem = AddressOperand::parse(stream)?;
91            stream.expect_complete()?;
92            stream.expect(&PtxToken::Comma)?;
93            let a_desc = GeneralOperand::parse(stream)?;
94            stream.expect_complete()?;
95            stream.expect(&PtxToken::Comma)?;
96            let b_desc = GeneralOperand::parse(stream)?;
97            stream.expect_complete()?;
98            stream.expect(&PtxToken::Comma)?;
99            let idesc = GeneralOperand::parse(stream)?;
100            stream.expect_complete()?;
101            stream.expect(&PtxToken::Comma)?;
102            let enable_input_d = GeneralOperand::parse(stream)?;
103            stream.expect_complete()?;
104            let saved_pos = stream.position();
105            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
106            if !has_comma {
107                stream.set_position(saved_pos);
108            }
109            let saved_pos = stream.position();
110            let zero_column_mask_desc = match GeneralOperand::parse(stream) {
111                Ok(val) => Some(val),
112                Err(_) => {
113                    stream.set_position(saved_pos);
114                    None
115                }
116            };
117            stream.expect_complete()?;
118            stream.expect_complete()?;
119            stream.expect(&PtxToken::Semicolon)?;
120            Ok(Tcgen05MmaWsCtaGroup1KindCollectorUsage {
121                mma,
122                ws,
123                cta_group_1,
124                kind,
125                collector_usage,
126                d_tmem,
127                a_desc,
128                b_desc,
129                idesc,
130                enable_input_d,
131                zero_column_mask_desc,
132            })
133        }
134    }
135
136
137    impl PtxParser for Tcgen05MmaWsCtaGroup1KindCollectorUsage1 {
138        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
139            stream.expect_string("tcgen05")?;
140            stream.expect_string(".mma")?;
141            let mma = ();
142            stream.expect_complete()?;
143            stream.expect_string(".ws")?;
144            let ws = ();
145            stream.expect_complete()?;
146            stream.expect_string(".cta_group::1")?;
147            let cta_group_1 = ();
148            stream.expect_complete()?;
149            let kind = Kind::parse(stream)?;
150            stream.expect_complete()?;
151            let saved_pos = stream.position();
152            let collector_usage = stream.expect_string(".collector_usage").is_ok();
153            if !collector_usage {
154                stream.set_position(saved_pos);
155            }
156            stream.expect_complete()?;
157            let d_tmem = AddressOperand::parse(stream)?;
158            stream.expect_complete()?;
159            stream.expect(&PtxToken::Comma)?;
160            let a_tmem = AddressOperand::parse(stream)?;
161            stream.expect_complete()?;
162            stream.expect(&PtxToken::Comma)?;
163            let b_desc = GeneralOperand::parse(stream)?;
164            stream.expect_complete()?;
165            stream.expect(&PtxToken::Comma)?;
166            let idesc = GeneralOperand::parse(stream)?;
167            stream.expect_complete()?;
168            stream.expect(&PtxToken::Comma)?;
169            let enable_input_d = GeneralOperand::parse(stream)?;
170            stream.expect_complete()?;
171            let saved_pos = stream.position();
172            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
173            if !has_comma {
174                stream.set_position(saved_pos);
175            }
176            let saved_pos = stream.position();
177            let zero_column_mask_desc = match GeneralOperand::parse(stream) {
178                Ok(val) => Some(val),
179                Err(_) => {
180                    stream.set_position(saved_pos);
181                    None
182                }
183            };
184            stream.expect_complete()?;
185            stream.expect_complete()?;
186            stream.expect(&PtxToken::Semicolon)?;
187            Ok(Tcgen05MmaWsCtaGroup1KindCollectorUsage1 {
188                mma,
189                ws,
190                cta_group_1,
191                kind,
192                collector_usage,
193                d_tmem,
194                a_tmem,
195                b_desc,
196                idesc,
197                enable_input_d,
198                zero_column_mask_desc,
199            })
200        }
201    }
202
203
204}
205
206pub mod section_1 {
207    use super::*;
208    use crate::r#type::instruction::tcgen05_mma_ws::section_1::*;
209
210    // ============================================================================
211    // Generated enum parsers
212    // ============================================================================
213
214    impl PtxParser for Buffer {
215        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
216            // Try B0
217            {
218                let saved_pos = stream.position();
219                if stream.expect_string("::b0").is_ok() {
220                    return Ok(Buffer::B0);
221                }
222                stream.set_position(saved_pos);
223            }
224            let saved_pos = stream.position();
225            // Try B1
226            {
227                let saved_pos = stream.position();
228                if stream.expect_string("::b1").is_ok() {
229                    return Ok(Buffer::B1);
230                }
231                stream.set_position(saved_pos);
232            }
233            stream.set_position(saved_pos);
234            let saved_pos = stream.position();
235            // Try B2
236            {
237                let saved_pos = stream.position();
238                if stream.expect_string("::b2").is_ok() {
239                    return Ok(Buffer::B2);
240                }
241                stream.set_position(saved_pos);
242            }
243            stream.set_position(saved_pos);
244            let saved_pos = stream.position();
245            // Try B3
246            {
247                let saved_pos = stream.position();
248                if stream.expect_string("::b3").is_ok() {
249                    return Ok(Buffer::B3);
250                }
251                stream.set_position(saved_pos);
252            }
253            stream.set_position(saved_pos);
254            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
255            let expected = &["::b0", "::b1", "::b2", "::b3"];
256            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
257            Err(crate::parser::unexpected_value(span, expected, found))
258        }
259    }
260
261    impl PtxParser for CollectorUsage {
262        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
263            // Try CollectorBufferOp
264            {
265                let saved_seq_pos = stream.position();
266                match (|| -> Result<_, PtxParseError> {
267            stream.expect_string(".collector")?;
268            let collector = ();
269            let buffer = Buffer::parse(stream)?;
270            let op = Op::parse(stream)?;
271                    Ok((collector, buffer, op))
272                })() {
273                    Ok((collector, buffer, op)) => {
274                        return Ok(CollectorUsage::CollectorBufferOp(collector, buffer, op));
275                    }
276                    Err(_) => {
277                        stream.set_position(saved_seq_pos);
278                    }
279                }
280            }
281            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
282            let expected = &["<complex>"];
283            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
284            Err(crate::parser::unexpected_value(span, expected, found))
285        }
286    }
287
288    impl PtxParser for Op {
289        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
290            // Try Lastuse
291            {
292                let saved_pos = stream.position();
293                if stream.expect_string("::lastuse").is_ok() {
294                    return Ok(Op::Lastuse);
295                }
296                stream.set_position(saved_pos);
297            }
298            let saved_pos = stream.position();
299            // Try Discard
300            {
301                let saved_pos = stream.position();
302                if stream.expect_string("::discard").is_ok() {
303                    return Ok(Op::Discard);
304                }
305                stream.set_position(saved_pos);
306            }
307            stream.set_position(saved_pos);
308            let saved_pos = stream.position();
309            // Try Fill
310            {
311                let saved_pos = stream.position();
312                if stream.expect_string("::fill").is_ok() {
313                    return Ok(Op::Fill);
314                }
315                stream.set_position(saved_pos);
316            }
317            stream.set_position(saved_pos);
318            let saved_pos = stream.position();
319            // Try Use
320            {
321                let saved_pos = stream.position();
322                if stream.expect_string("::use").is_ok() {
323                    return Ok(Op::Use);
324                }
325                stream.set_position(saved_pos);
326            }
327            stream.set_position(saved_pos);
328            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
329            let expected = &["::lastuse", "::discard", "::fill", "::use"];
330            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
331            Err(crate::parser::unexpected_value(span, expected, found))
332        }
333    }
334
335    impl PtxParser for Tcgen05MmaWsCtaGroup1KindI8CollectorUsage {
336        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
337            stream.expect_string("tcgen05")?;
338            stream.expect_string(".mma")?;
339            let mma = ();
340            stream.expect_complete()?;
341            stream.expect_string(".ws")?;
342            let ws = ();
343            stream.expect_complete()?;
344            stream.expect_string(".cta_group::1")?;
345            let cta_group_1 = ();
346            stream.expect_complete()?;
347            stream.expect_string(".kind::i8")?;
348            let kind_i8 = ();
349            stream.expect_complete()?;
350            let saved_pos = stream.position();
351            let collector_usage = match CollectorUsage::parse(stream) {
352                Ok(val) => Some(val),
353                Err(_) => {
354                    stream.set_position(saved_pos);
355                    None
356                }
357            };
358            stream.expect_complete()?;
359            let d_tmem = AddressOperand::parse(stream)?;
360            stream.expect_complete()?;
361            stream.expect(&PtxToken::Comma)?;
362            let a_desc = GeneralOperand::parse(stream)?;
363            stream.expect_complete()?;
364            stream.expect(&PtxToken::Comma)?;
365            let b_desc = GeneralOperand::parse(stream)?;
366            stream.expect_complete()?;
367            stream.expect(&PtxToken::Comma)?;
368            let idesc = GeneralOperand::parse(stream)?;
369            stream.expect_complete()?;
370            stream.expect(&PtxToken::Comma)?;
371            let enable_input_d = GeneralOperand::parse(stream)?;
372            stream.expect_complete()?;
373            let saved_pos = stream.position();
374            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
375            if !has_comma {
376                stream.set_position(saved_pos);
377            }
378            let saved_pos = stream.position();
379            let zero_column_mask_desc = match GeneralOperand::parse(stream) {
380                Ok(val) => Some(val),
381                Err(_) => {
382                    stream.set_position(saved_pos);
383                    None
384                }
385            };
386            stream.expect_complete()?;
387            stream.expect_complete()?;
388            stream.expect(&PtxToken::Semicolon)?;
389            Ok(Tcgen05MmaWsCtaGroup1KindI8CollectorUsage {
390                mma,
391                ws,
392                cta_group_1,
393                kind_i8,
394                collector_usage,
395                d_tmem,
396                a_desc,
397                b_desc,
398                idesc,
399                enable_input_d,
400                zero_column_mask_desc,
401            })
402        }
403    }
404
405
406    impl PtxParser for Tcgen05MmaWsCtaGroup1KindI8CollectorUsage1 {
407        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
408            stream.expect_string("tcgen05")?;
409            stream.expect_string(".mma")?;
410            let mma = ();
411            stream.expect_complete()?;
412            stream.expect_string(".ws")?;
413            let ws = ();
414            stream.expect_complete()?;
415            stream.expect_string(".cta_group::1")?;
416            let cta_group_1 = ();
417            stream.expect_complete()?;
418            stream.expect_string(".kind::i8")?;
419            let kind_i8 = ();
420            stream.expect_complete()?;
421            let saved_pos = stream.position();
422            let collector_usage = match CollectorUsage::parse(stream) {
423                Ok(val) => Some(val),
424                Err(_) => {
425                    stream.set_position(saved_pos);
426                    None
427                }
428            };
429            stream.expect_complete()?;
430            let d_tmem = AddressOperand::parse(stream)?;
431            stream.expect_complete()?;
432            stream.expect(&PtxToken::Comma)?;
433            let a_tmem = AddressOperand::parse(stream)?;
434            stream.expect_complete()?;
435            stream.expect(&PtxToken::Comma)?;
436            let b_desc = GeneralOperand::parse(stream)?;
437            stream.expect_complete()?;
438            stream.expect(&PtxToken::Comma)?;
439            let idesc = GeneralOperand::parse(stream)?;
440            stream.expect_complete()?;
441            stream.expect(&PtxToken::Comma)?;
442            let enable_input_d = GeneralOperand::parse(stream)?;
443            stream.expect_complete()?;
444            let saved_pos = stream.position();
445            let has_comma = stream.expect(&PtxToken::Comma).is_ok();
446            if !has_comma {
447                stream.set_position(saved_pos);
448            }
449            let saved_pos = stream.position();
450            let zero_column_mask_desc = match GeneralOperand::parse(stream) {
451                Ok(val) => Some(val),
452                Err(_) => {
453                    stream.set_position(saved_pos);
454                    None
455                }
456            };
457            stream.expect_complete()?;
458            stream.expect_complete()?;
459            stream.expect(&PtxToken::Semicolon)?;
460            Ok(Tcgen05MmaWsCtaGroup1KindI8CollectorUsage1 {
461                mma,
462                ws,
463                cta_group_1,
464                kind_i8,
465                collector_usage,
466                d_tmem,
467                a_tmem,
468                b_desc,
469                idesc,
470                enable_input_d,
471                zero_column_mask_desc,
472            })
473        }
474    }
475
476
477}
478
ptx_parser/parser/instruction/tcgen05_mma_ws.rs

ptx_parser/parser/instruction/
tcgen05_mma_ws.rs
