ptx_parser/parser/instruction/
tcgen05_ld.rs

1//! Original PTX specification:
2//!
3//! // Base load instruction:
4//! tcgen05.ld.sync.aligned.shape1.num{.pack}.b32    r, [taddr];
5//! tcgen05.ld.sync.aligned.shape2.num{.pack}.b32    r, [taddr], immHalfSplitoff;
6//! .shape1 = { .16x64b, .16x128b, .16x256b, .32x32b };
7//! .shape2 = { .16x32bx2 };
8//! .num    = { .x1, .x2, .x4, .x8, .x16, .x32, .x64, .x128 };
9//! .pack   = { .pack::16b };
10//! // Floating point type load along with reduction :
11//! tcgen05.ld.red.sync.aligned.shape3.num.redOp{.abs}{.NaN}.f32 r, redval, [taddr];
12//! tcgen05.ld.red.sync.aligned.shape4.num.redOp{.abs}{.NaN}.f32 r, redval, [taddr], immHalfSplitoff;
13//! // Integer type load along with reduction :
14//! tcgen05.ld.red.sync.aligned.shape3.num.redOp.type r, redval, [taddr];
15//! tcgen05.ld.red.sync.aligned.shape4.num.redOp.type r, redval, [taddr], immHalfSplitoff;
16//! .shape3 = { .32x32b   };
17//! .shape4 = { .16x32bx2 };
18//! .redOp  = { .min, .max };
19//! .type   = { .u32, .s32 };
20
21#![allow(unused)]
22
23use crate::lexer::PtxToken;
24use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
25use crate::r#type::common::*;
26
27pub mod section_0 {
28    use super::*;
29    use crate::r#type::instruction::tcgen05_ld::section_0::*;
30
31    // ============================================================================
32    // Generated enum parsers
33    // ============================================================================
34
35    impl PtxParser for Num {
36        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
37            // Try X128
38            {
39                let saved_pos = stream.position();
40                if stream.expect_string(".x128").is_ok() {
41                    return Ok(Num::X128);
42                }
43                stream.set_position(saved_pos);
44            }
45            let saved_pos = stream.position();
46            // Try X16
47            {
48                let saved_pos = stream.position();
49                if stream.expect_string(".x16").is_ok() {
50                    return Ok(Num::X16);
51                }
52                stream.set_position(saved_pos);
53            }
54            stream.set_position(saved_pos);
55            let saved_pos = stream.position();
56            // Try X32
57            {
58                let saved_pos = stream.position();
59                if stream.expect_string(".x32").is_ok() {
60                    return Ok(Num::X32);
61                }
62                stream.set_position(saved_pos);
63            }
64            stream.set_position(saved_pos);
65            let saved_pos = stream.position();
66            // Try X64
67            {
68                let saved_pos = stream.position();
69                if stream.expect_string(".x64").is_ok() {
70                    return Ok(Num::X64);
71                }
72                stream.set_position(saved_pos);
73            }
74            stream.set_position(saved_pos);
75            let saved_pos = stream.position();
76            // Try X1
77            {
78                let saved_pos = stream.position();
79                if stream.expect_string(".x1").is_ok() {
80                    return Ok(Num::X1);
81                }
82                stream.set_position(saved_pos);
83            }
84            stream.set_position(saved_pos);
85            let saved_pos = stream.position();
86            // Try X2
87            {
88                let saved_pos = stream.position();
89                if stream.expect_string(".x2").is_ok() {
90                    return Ok(Num::X2);
91                }
92                stream.set_position(saved_pos);
93            }
94            stream.set_position(saved_pos);
95            let saved_pos = stream.position();
96            // Try X4
97            {
98                let saved_pos = stream.position();
99                if stream.expect_string(".x4").is_ok() {
100                    return Ok(Num::X4);
101                }
102                stream.set_position(saved_pos);
103            }
104            stream.set_position(saved_pos);
105            let saved_pos = stream.position();
106            // Try X8
107            {
108                let saved_pos = stream.position();
109                if stream.expect_string(".x8").is_ok() {
110                    return Ok(Num::X8);
111                }
112                stream.set_position(saved_pos);
113            }
114            stream.set_position(saved_pos);
115            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
116            let expected = &[".x128", ".x16", ".x32", ".x64", ".x1", ".x2", ".x4", ".x8"];
117            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
118            Err(crate::parser::unexpected_value(span, expected, found))
119        }
120    }
121
122    impl PtxParser for Pack {
123        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
124            // Try Pack16b
125            {
126                let saved_pos = stream.position();
127                if stream.expect_string(".pack::16b").is_ok() {
128                    return Ok(Pack::Pack16b);
129                }
130                stream.set_position(saved_pos);
131            }
132            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
133            let expected = &[".pack::16b"];
134            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
135            Err(crate::parser::unexpected_value(span, expected, found))
136        }
137    }
138
139    impl PtxParser for Redop {
140        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
141            // Try Min
142            {
143                let saved_pos = stream.position();
144                if stream.expect_string(".min").is_ok() {
145                    return Ok(Redop::Min);
146                }
147                stream.set_position(saved_pos);
148            }
149            let saved_pos = stream.position();
150            // Try Max
151            {
152                let saved_pos = stream.position();
153                if stream.expect_string(".max").is_ok() {
154                    return Ok(Redop::Max);
155                }
156                stream.set_position(saved_pos);
157            }
158            stream.set_position(saved_pos);
159            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
160            let expected = &[".min", ".max"];
161            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
162            Err(crate::parser::unexpected_value(span, expected, found))
163        }
164    }
165
166    impl PtxParser for Shape1 {
167        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
168            // Try _16x128b
169            {
170                let saved_pos = stream.position();
171                if stream.expect_string(".16x128b").is_ok() {
172                    return Ok(Shape1::_16x128b);
173                }
174                stream.set_position(saved_pos);
175            }
176            let saved_pos = stream.position();
177            // Try _16x256b
178            {
179                let saved_pos = stream.position();
180                if stream.expect_string(".16x256b").is_ok() {
181                    return Ok(Shape1::_16x256b);
182                }
183                stream.set_position(saved_pos);
184            }
185            stream.set_position(saved_pos);
186            let saved_pos = stream.position();
187            // Try _16x64b
188            {
189                let saved_pos = stream.position();
190                if stream.expect_string(".16x64b").is_ok() {
191                    return Ok(Shape1::_16x64b);
192                }
193                stream.set_position(saved_pos);
194            }
195            stream.set_position(saved_pos);
196            let saved_pos = stream.position();
197            // Try _32x32b
198            {
199                let saved_pos = stream.position();
200                if stream.expect_string(".32x32b").is_ok() {
201                    return Ok(Shape1::_32x32b);
202                }
203                stream.set_position(saved_pos);
204            }
205            stream.set_position(saved_pos);
206            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
207            let expected = &[".16x128b", ".16x256b", ".16x64b", ".32x32b"];
208            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
209            Err(crate::parser::unexpected_value(span, expected, found))
210        }
211    }
212
213    impl PtxParser for Shape2 {
214        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
215            // Try _16x32bx2
216            {
217                let saved_pos = stream.position();
218                if stream.expect_string(".16x32bx2").is_ok() {
219                    return Ok(Shape2::_16x32bx2);
220                }
221                stream.set_position(saved_pos);
222            }
223            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
224            let expected = &[".16x32bx2"];
225            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
226            Err(crate::parser::unexpected_value(span, expected, found))
227        }
228    }
229
230    impl PtxParser for Shape3 {
231        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
232            // Try _32x32b
233            {
234                let saved_pos = stream.position();
235                if stream.expect_string(".32x32b").is_ok() {
236                    return Ok(Shape3::_32x32b);
237                }
238                stream.set_position(saved_pos);
239            }
240            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
241            let expected = &[".32x32b"];
242            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
243            Err(crate::parser::unexpected_value(span, expected, found))
244        }
245    }
246
247    impl PtxParser for Shape4 {
248        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
249            // Try _16x32bx2
250            {
251                let saved_pos = stream.position();
252                if stream.expect_string(".16x32bx2").is_ok() {
253                    return Ok(Shape4::_16x32bx2);
254                }
255                stream.set_position(saved_pos);
256            }
257            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
258            let expected = &[".16x32bx2"];
259            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
260            Err(crate::parser::unexpected_value(span, expected, found))
261        }
262    }
263
264    impl PtxParser for Type {
265        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
266            // Try U32
267            {
268                let saved_pos = stream.position();
269                if stream.expect_string(".u32").is_ok() {
270                    return Ok(Type::U32);
271                }
272                stream.set_position(saved_pos);
273            }
274            let saved_pos = stream.position();
275            // Try S32
276            {
277                let saved_pos = stream.position();
278                if stream.expect_string(".s32").is_ok() {
279                    return Ok(Type::S32);
280                }
281                stream.set_position(saved_pos);
282            }
283            stream.set_position(saved_pos);
284            let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
285            let expected = &[".u32", ".s32"];
286            let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
287            Err(crate::parser::unexpected_value(span, expected, found))
288        }
289    }
290
291    impl PtxParser for Tcgen05LdSyncAlignedShape1NumPackB32 {
292        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
293            stream.expect_string("tcgen05")?;
294            stream.expect_string(".ld")?;
295            let ld = ();
296            stream.expect_complete()?;
297            stream.expect_string(".sync")?;
298            let sync = ();
299            stream.expect_complete()?;
300            stream.expect_string(".aligned")?;
301            let aligned = ();
302            stream.expect_complete()?;
303            let shape1 = Shape1::parse(stream)?;
304            stream.expect_complete()?;
305            let num = Num::parse(stream)?;
306            stream.expect_complete()?;
307            let saved_pos = stream.position();
308            let pack = match Pack::parse(stream) {
309                Ok(val) => Some(val),
310                Err(_) => {
311                    stream.set_position(saved_pos);
312                    None
313                }
314            };
315            stream.expect_complete()?;
316            stream.expect_string(".b32")?;
317            let b32 = ();
318            stream.expect_complete()?;
319            let r = GeneralOperand::parse(stream)?;
320            stream.expect_complete()?;
321            stream.expect(&PtxToken::Comma)?;
322            let taddr = AddressOperand::parse(stream)?;
323            stream.expect_complete()?;
324            stream.expect_complete()?;
325            stream.expect(&PtxToken::Semicolon)?;
326            Ok(Tcgen05LdSyncAlignedShape1NumPackB32 {
327                ld,
328                sync,
329                aligned,
330                shape1,
331                num,
332                pack,
333                b32,
334                r,
335                taddr,
336            })
337        }
338    }
339
340
341    impl PtxParser for Tcgen05LdSyncAlignedShape2NumPackB32 {
342        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
343            stream.expect_string("tcgen05")?;
344            stream.expect_string(".ld")?;
345            let ld = ();
346            stream.expect_complete()?;
347            stream.expect_string(".sync")?;
348            let sync = ();
349            stream.expect_complete()?;
350            stream.expect_string(".aligned")?;
351            let aligned = ();
352            stream.expect_complete()?;
353            let shape2 = Shape2::parse(stream)?;
354            stream.expect_complete()?;
355            let num = Num::parse(stream)?;
356            stream.expect_complete()?;
357            let saved_pos = stream.position();
358            let pack = match Pack::parse(stream) {
359                Ok(val) => Some(val),
360                Err(_) => {
361                    stream.set_position(saved_pos);
362                    None
363                }
364            };
365            stream.expect_complete()?;
366            stream.expect_string(".b32")?;
367            let b32 = ();
368            stream.expect_complete()?;
369            let r = GeneralOperand::parse(stream)?;
370            stream.expect_complete()?;
371            stream.expect(&PtxToken::Comma)?;
372            let taddr = AddressOperand::parse(stream)?;
373            stream.expect_complete()?;
374            stream.expect(&PtxToken::Comma)?;
375            let immhalfsplitoff = GeneralOperand::parse(stream)?;
376            stream.expect_complete()?;
377            stream.expect_complete()?;
378            stream.expect(&PtxToken::Semicolon)?;
379            Ok(Tcgen05LdSyncAlignedShape2NumPackB32 {
380                ld,
381                sync,
382                aligned,
383                shape2,
384                num,
385                pack,
386                b32,
387                r,
388                taddr,
389                immhalfsplitoff,
390            })
391        }
392    }
393
394
395    impl PtxParser for Tcgen05LdRedSyncAlignedShape3NumRedopAbsNanF32 {
396        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
397            stream.expect_string("tcgen05")?;
398            stream.expect_string(".ld")?;
399            let ld = ();
400            stream.expect_complete()?;
401            stream.expect_string(".red")?;
402            let red = ();
403            stream.expect_complete()?;
404            stream.expect_string(".sync")?;
405            let sync = ();
406            stream.expect_complete()?;
407            stream.expect_string(".aligned")?;
408            let aligned = ();
409            stream.expect_complete()?;
410            let shape3 = Shape3::parse(stream)?;
411            stream.expect_complete()?;
412            let num = Num::parse(stream)?;
413            stream.expect_complete()?;
414            let redop = Redop::parse(stream)?;
415            stream.expect_complete()?;
416            let saved_pos = stream.position();
417            let abs = stream.expect_string(".abs").is_ok();
418            if !abs {
419                stream.set_position(saved_pos);
420            }
421            stream.expect_complete()?;
422            let saved_pos = stream.position();
423            let nan = stream.expect_string(".NaN").is_ok();
424            if !nan {
425                stream.set_position(saved_pos);
426            }
427            stream.expect_complete()?;
428            stream.expect_string(".f32")?;
429            let f32 = ();
430            stream.expect_complete()?;
431            let r = GeneralOperand::parse(stream)?;
432            stream.expect_complete()?;
433            stream.expect(&PtxToken::Comma)?;
434            let redval = GeneralOperand::parse(stream)?;
435            stream.expect_complete()?;
436            stream.expect(&PtxToken::Comma)?;
437            let taddr = AddressOperand::parse(stream)?;
438            stream.expect_complete()?;
439            stream.expect_complete()?;
440            stream.expect(&PtxToken::Semicolon)?;
441            Ok(Tcgen05LdRedSyncAlignedShape3NumRedopAbsNanF32 {
442                ld,
443                red,
444                sync,
445                aligned,
446                shape3,
447                num,
448                redop,
449                abs,
450                nan,
451                f32,
452                r,
453                redval,
454                taddr,
455            })
456        }
457    }
458
459
460    impl PtxParser for Tcgen05LdRedSyncAlignedShape4NumRedopAbsNanF32 {
461        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
462            stream.expect_string("tcgen05")?;
463            stream.expect_string(".ld")?;
464            let ld = ();
465            stream.expect_complete()?;
466            stream.expect_string(".red")?;
467            let red = ();
468            stream.expect_complete()?;
469            stream.expect_string(".sync")?;
470            let sync = ();
471            stream.expect_complete()?;
472            stream.expect_string(".aligned")?;
473            let aligned = ();
474            stream.expect_complete()?;
475            let shape4 = Shape4::parse(stream)?;
476            stream.expect_complete()?;
477            let num = Num::parse(stream)?;
478            stream.expect_complete()?;
479            let redop = Redop::parse(stream)?;
480            stream.expect_complete()?;
481            let saved_pos = stream.position();
482            let abs = stream.expect_string(".abs").is_ok();
483            if !abs {
484                stream.set_position(saved_pos);
485            }
486            stream.expect_complete()?;
487            let saved_pos = stream.position();
488            let nan = stream.expect_string(".NaN").is_ok();
489            if !nan {
490                stream.set_position(saved_pos);
491            }
492            stream.expect_complete()?;
493            stream.expect_string(".f32")?;
494            let f32 = ();
495            stream.expect_complete()?;
496            let r = GeneralOperand::parse(stream)?;
497            stream.expect_complete()?;
498            stream.expect(&PtxToken::Comma)?;
499            let redval = GeneralOperand::parse(stream)?;
500            stream.expect_complete()?;
501            stream.expect(&PtxToken::Comma)?;
502            let taddr = AddressOperand::parse(stream)?;
503            stream.expect_complete()?;
504            stream.expect(&PtxToken::Comma)?;
505            let immhalfsplitoff = GeneralOperand::parse(stream)?;
506            stream.expect_complete()?;
507            stream.expect_complete()?;
508            stream.expect(&PtxToken::Semicolon)?;
509            Ok(Tcgen05LdRedSyncAlignedShape4NumRedopAbsNanF32 {
510                ld,
511                red,
512                sync,
513                aligned,
514                shape4,
515                num,
516                redop,
517                abs,
518                nan,
519                f32,
520                r,
521                redval,
522                taddr,
523                immhalfsplitoff,
524            })
525        }
526    }
527
528
529    impl PtxParser for Tcgen05LdRedSyncAlignedShape3NumRedopType {
530        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
531            stream.expect_string("tcgen05")?;
532            stream.expect_string(".ld")?;
533            let ld = ();
534            stream.expect_complete()?;
535            stream.expect_string(".red")?;
536            let red = ();
537            stream.expect_complete()?;
538            stream.expect_string(".sync")?;
539            let sync = ();
540            stream.expect_complete()?;
541            stream.expect_string(".aligned")?;
542            let aligned = ();
543            stream.expect_complete()?;
544            let shape3 = Shape3::parse(stream)?;
545            stream.expect_complete()?;
546            let num = Num::parse(stream)?;
547            stream.expect_complete()?;
548            let redop = Redop::parse(stream)?;
549            stream.expect_complete()?;
550            let type_ = Type::parse(stream)?;
551            stream.expect_complete()?;
552            let r = GeneralOperand::parse(stream)?;
553            stream.expect_complete()?;
554            stream.expect(&PtxToken::Comma)?;
555            let redval = GeneralOperand::parse(stream)?;
556            stream.expect_complete()?;
557            stream.expect(&PtxToken::Comma)?;
558            let taddr = AddressOperand::parse(stream)?;
559            stream.expect_complete()?;
560            stream.expect_complete()?;
561            stream.expect(&PtxToken::Semicolon)?;
562            Ok(Tcgen05LdRedSyncAlignedShape3NumRedopType {
563                ld,
564                red,
565                sync,
566                aligned,
567                shape3,
568                num,
569                redop,
570                type_,
571                r,
572                redval,
573                taddr,
574            })
575        }
576    }
577
578
579    impl PtxParser for Tcgen05LdRedSyncAlignedShape4NumRedopType {
580        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
581            stream.expect_string("tcgen05")?;
582            stream.expect_string(".ld")?;
583            let ld = ();
584            stream.expect_complete()?;
585            stream.expect_string(".red")?;
586            let red = ();
587            stream.expect_complete()?;
588            stream.expect_string(".sync")?;
589            let sync = ();
590            stream.expect_complete()?;
591            stream.expect_string(".aligned")?;
592            let aligned = ();
593            stream.expect_complete()?;
594            let shape4 = Shape4::parse(stream)?;
595            stream.expect_complete()?;
596            let num = Num::parse(stream)?;
597            stream.expect_complete()?;
598            let redop = Redop::parse(stream)?;
599            stream.expect_complete()?;
600            let type_ = Type::parse(stream)?;
601            stream.expect_complete()?;
602            let r = GeneralOperand::parse(stream)?;
603            stream.expect_complete()?;
604            stream.expect(&PtxToken::Comma)?;
605            let redval = GeneralOperand::parse(stream)?;
606            stream.expect_complete()?;
607            stream.expect(&PtxToken::Comma)?;
608            let taddr = AddressOperand::parse(stream)?;
609            stream.expect_complete()?;
610            stream.expect(&PtxToken::Comma)?;
611            let immhalfsplitoff = GeneralOperand::parse(stream)?;
612            stream.expect_complete()?;
613            stream.expect_complete()?;
614            stream.expect(&PtxToken::Semicolon)?;
615            Ok(Tcgen05LdRedSyncAlignedShape4NumRedopType {
616                ld,
617                red,
618                sync,
619                aligned,
620                shape4,
621                num,
622                redop,
623                type_,
624                r,
625                redval,
626                taddr,
627                immhalfsplitoff,
628            })
629        }
630    }
631
632
633}
634
ptx_parser/parser/instruction/tcgen05_ld.rs

ptx_parser/parser/instruction/
tcgen05_ld.rs
