ptx_parser/parser/instruction/
tcgen05_mma_ws_sp.rs

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

ptx_parser/parser/instruction/
tcgen05_mma_ws_sp.rs
