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
116                .peek()
117                .map(|(_, s)| s.clone())
118                .unwrap_or(Span { start: 0, end: 0 });
119            let expected = &[".x128", ".x16", ".x32", ".x64", ".x1", ".x2", ".x4", ".x8"];
120            let found = stream
121                .peek()
122                .map(|(t, _)| format!("{:?}", t))
123                .unwrap_or_else(|_| "<end of input>".to_string());
124            Err(crate::parser::unexpected_value(span, expected, found))
125        }
126    }
127
128    impl PtxParser for Pack {
129        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
130            // Try Pack16b
131            {
132                let saved_pos = stream.position();
133                if stream.expect_string(".pack::16b").is_ok() {
134                    return Ok(Pack::Pack16b);
135                }
136                stream.set_position(saved_pos);
137            }
138            let span = stream
139                .peek()
140                .map(|(_, s)| s.clone())
141                .unwrap_or(Span { start: 0, end: 0 });
142            let expected = &[".pack::16b"];
143            let found = stream
144                .peek()
145                .map(|(t, _)| format!("{:?}", t))
146                .unwrap_or_else(|_| "<end of input>".to_string());
147            Err(crate::parser::unexpected_value(span, expected, found))
148        }
149    }
150
151    impl PtxParser for Redop {
152        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
153            // Try Min
154            {
155                let saved_pos = stream.position();
156                if stream.expect_string(".min").is_ok() {
157                    return Ok(Redop::Min);
158                }
159                stream.set_position(saved_pos);
160            }
161            let saved_pos = stream.position();
162            // Try Max
163            {
164                let saved_pos = stream.position();
165                if stream.expect_string(".max").is_ok() {
166                    return Ok(Redop::Max);
167                }
168                stream.set_position(saved_pos);
169            }
170            stream.set_position(saved_pos);
171            let span = stream
172                .peek()
173                .map(|(_, s)| s.clone())
174                .unwrap_or(Span { start: 0, end: 0 });
175            let expected = &[".min", ".max"];
176            let found = stream
177                .peek()
178                .map(|(t, _)| format!("{:?}", t))
179                .unwrap_or_else(|_| "<end of input>".to_string());
180            Err(crate::parser::unexpected_value(span, expected, found))
181        }
182    }
183
184    impl PtxParser for Shape1 {
185        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
186            // Try _16x128b
187            {
188                let saved_pos = stream.position();
189                if stream.expect_string(".16x128b").is_ok() {
190                    return Ok(Shape1::_16x128b);
191                }
192                stream.set_position(saved_pos);
193            }
194            let saved_pos = stream.position();
195            // Try _16x256b
196            {
197                let saved_pos = stream.position();
198                if stream.expect_string(".16x256b").is_ok() {
199                    return Ok(Shape1::_16x256b);
200                }
201                stream.set_position(saved_pos);
202            }
203            stream.set_position(saved_pos);
204            let saved_pos = stream.position();
205            // Try _16x64b
206            {
207                let saved_pos = stream.position();
208                if stream.expect_string(".16x64b").is_ok() {
209                    return Ok(Shape1::_16x64b);
210                }
211                stream.set_position(saved_pos);
212            }
213            stream.set_position(saved_pos);
214            let saved_pos = stream.position();
215            // Try _32x32b
216            {
217                let saved_pos = stream.position();
218                if stream.expect_string(".32x32b").is_ok() {
219                    return Ok(Shape1::_32x32b);
220                }
221                stream.set_position(saved_pos);
222            }
223            stream.set_position(saved_pos);
224            let span = stream
225                .peek()
226                .map(|(_, s)| s.clone())
227                .unwrap_or(Span { start: 0, end: 0 });
228            let expected = &[".16x128b", ".16x256b", ".16x64b", ".32x32b"];
229            let found = stream
230                .peek()
231                .map(|(t, _)| format!("{:?}", t))
232                .unwrap_or_else(|_| "<end of input>".to_string());
233            Err(crate::parser::unexpected_value(span, expected, found))
234        }
235    }
236
237    impl PtxParser for Shape2 {
238        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
239            // Try _16x32bx2
240            {
241                let saved_pos = stream.position();
242                if stream.expect_string(".16x32bx2").is_ok() {
243                    return Ok(Shape2::_16x32bx2);
244                }
245                stream.set_position(saved_pos);
246            }
247            let span = stream
248                .peek()
249                .map(|(_, s)| s.clone())
250                .unwrap_or(Span { start: 0, end: 0 });
251            let expected = &[".16x32bx2"];
252            let found = stream
253                .peek()
254                .map(|(t, _)| format!("{:?}", t))
255                .unwrap_or_else(|_| "<end of input>".to_string());
256            Err(crate::parser::unexpected_value(span, expected, found))
257        }
258    }
259
260    impl PtxParser for Shape3 {
261        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
262            // Try _32x32b
263            {
264                let saved_pos = stream.position();
265                if stream.expect_string(".32x32b").is_ok() {
266                    return Ok(Shape3::_32x32b);
267                }
268                stream.set_position(saved_pos);
269            }
270            let span = stream
271                .peek()
272                .map(|(_, s)| s.clone())
273                .unwrap_or(Span { start: 0, end: 0 });
274            let expected = &[".32x32b"];
275            let found = stream
276                .peek()
277                .map(|(t, _)| format!("{:?}", t))
278                .unwrap_or_else(|_| "<end of input>".to_string());
279            Err(crate::parser::unexpected_value(span, expected, found))
280        }
281    }
282
283    impl PtxParser for Shape4 {
284        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
285            // Try _16x32bx2
286            {
287                let saved_pos = stream.position();
288                if stream.expect_string(".16x32bx2").is_ok() {
289                    return Ok(Shape4::_16x32bx2);
290                }
291                stream.set_position(saved_pos);
292            }
293            let span = stream
294                .peek()
295                .map(|(_, s)| s.clone())
296                .unwrap_or(Span { start: 0, end: 0 });
297            let expected = &[".16x32bx2"];
298            let found = stream
299                .peek()
300                .map(|(t, _)| format!("{:?}", t))
301                .unwrap_or_else(|_| "<end of input>".to_string());
302            Err(crate::parser::unexpected_value(span, expected, found))
303        }
304    }
305
306    impl PtxParser for Type {
307        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
308            // Try U32
309            {
310                let saved_pos = stream.position();
311                if stream.expect_string(".u32").is_ok() {
312                    return Ok(Type::U32);
313                }
314                stream.set_position(saved_pos);
315            }
316            let saved_pos = stream.position();
317            // Try S32
318            {
319                let saved_pos = stream.position();
320                if stream.expect_string(".s32").is_ok() {
321                    return Ok(Type::S32);
322                }
323                stream.set_position(saved_pos);
324            }
325            stream.set_position(saved_pos);
326            let span = stream
327                .peek()
328                .map(|(_, s)| s.clone())
329                .unwrap_or(Span { start: 0, end: 0 });
330            let expected = &[".u32", ".s32"];
331            let found = stream
332                .peek()
333                .map(|(t, _)| format!("{:?}", t))
334                .unwrap_or_else(|_| "<end of input>".to_string());
335            Err(crate::parser::unexpected_value(span, expected, found))
336        }
337    }
338
339    impl PtxParser for Tcgen05LdSyncAlignedShape1NumPackB32 {
340        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
341            stream.expect_string("tcgen05")?;
342            stream.expect_string(".ld")?;
343            let ld = ();
344            stream.expect_complete()?;
345            stream.expect_string(".sync")?;
346            let sync = ();
347            stream.expect_complete()?;
348            stream.expect_string(".aligned")?;
349            let aligned = ();
350            stream.expect_complete()?;
351            let shape1 = Shape1::parse(stream)?;
352            stream.expect_complete()?;
353            let num = Num::parse(stream)?;
354            stream.expect_complete()?;
355            let saved_pos = stream.position();
356            let pack = match Pack::parse(stream) {
357                Ok(val) => Some(val),
358                Err(_) => {
359                    stream.set_position(saved_pos);
360                    None
361                }
362            };
363            stream.expect_complete()?;
364            stream.expect_string(".b32")?;
365            let b32 = ();
366            stream.expect_complete()?;
367            let r = GeneralOperand::parse(stream)?;
368            stream.expect_complete()?;
369            stream.expect(&PtxToken::Comma)?;
370            let taddr = AddressOperand::parse(stream)?;
371            stream.expect_complete()?;
372            stream.expect_complete()?;
373            stream.expect(&PtxToken::Semicolon)?;
374            Ok(Tcgen05LdSyncAlignedShape1NumPackB32 {
375                ld,
376                sync,
377                aligned,
378                shape1,
379                num,
380                pack,
381                b32,
382                r,
383                taddr,
384            })
385        }
386    }
387
388    impl PtxParser for Tcgen05LdSyncAlignedShape2NumPackB32 {
389        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
390            stream.expect_string("tcgen05")?;
391            stream.expect_string(".ld")?;
392            let ld = ();
393            stream.expect_complete()?;
394            stream.expect_string(".sync")?;
395            let sync = ();
396            stream.expect_complete()?;
397            stream.expect_string(".aligned")?;
398            let aligned = ();
399            stream.expect_complete()?;
400            let shape2 = Shape2::parse(stream)?;
401            stream.expect_complete()?;
402            let num = Num::parse(stream)?;
403            stream.expect_complete()?;
404            let saved_pos = stream.position();
405            let pack = match Pack::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_string(".b32")?;
414            let b32 = ();
415            stream.expect_complete()?;
416            let r = GeneralOperand::parse(stream)?;
417            stream.expect_complete()?;
418            stream.expect(&PtxToken::Comma)?;
419            let taddr = AddressOperand::parse(stream)?;
420            stream.expect_complete()?;
421            stream.expect(&PtxToken::Comma)?;
422            let immhalfsplitoff = GeneralOperand::parse(stream)?;
423            stream.expect_complete()?;
424            stream.expect_complete()?;
425            stream.expect(&PtxToken::Semicolon)?;
426            Ok(Tcgen05LdSyncAlignedShape2NumPackB32 {
427                ld,
428                sync,
429                aligned,
430                shape2,
431                num,
432                pack,
433                b32,
434                r,
435                taddr,
436                immhalfsplitoff,
437            })
438        }
439    }
440
441    impl PtxParser for Tcgen05LdRedSyncAlignedShape3NumRedopAbsNanF32 {
442        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
443            stream.expect_string("tcgen05")?;
444            stream.expect_string(".ld")?;
445            let ld = ();
446            stream.expect_complete()?;
447            stream.expect_string(".red")?;
448            let red = ();
449            stream.expect_complete()?;
450            stream.expect_string(".sync")?;
451            let sync = ();
452            stream.expect_complete()?;
453            stream.expect_string(".aligned")?;
454            let aligned = ();
455            stream.expect_complete()?;
456            let shape3 = Shape3::parse(stream)?;
457            stream.expect_complete()?;
458            let num = Num::parse(stream)?;
459            stream.expect_complete()?;
460            let redop = Redop::parse(stream)?;
461            stream.expect_complete()?;
462            let saved_pos = stream.position();
463            let abs = stream.expect_string(".abs").is_ok();
464            if !abs {
465                stream.set_position(saved_pos);
466            }
467            stream.expect_complete()?;
468            let saved_pos = stream.position();
469            let nan = stream.expect_string(".NaN").is_ok();
470            if !nan {
471                stream.set_position(saved_pos);
472            }
473            stream.expect_complete()?;
474            stream.expect_string(".f32")?;
475            let f32 = ();
476            stream.expect_complete()?;
477            let r = GeneralOperand::parse(stream)?;
478            stream.expect_complete()?;
479            stream.expect(&PtxToken::Comma)?;
480            let redval = GeneralOperand::parse(stream)?;
481            stream.expect_complete()?;
482            stream.expect(&PtxToken::Comma)?;
483            let taddr = AddressOperand::parse(stream)?;
484            stream.expect_complete()?;
485            stream.expect_complete()?;
486            stream.expect(&PtxToken::Semicolon)?;
487            Ok(Tcgen05LdRedSyncAlignedShape3NumRedopAbsNanF32 {
488                ld,
489                red,
490                sync,
491                aligned,
492                shape3,
493                num,
494                redop,
495                abs,
496                nan,
497                f32,
498                r,
499                redval,
500                taddr,
501            })
502        }
503    }
504
505    impl PtxParser for Tcgen05LdRedSyncAlignedShape4NumRedopAbsNanF32 {
506        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
507            stream.expect_string("tcgen05")?;
508            stream.expect_string(".ld")?;
509            let ld = ();
510            stream.expect_complete()?;
511            stream.expect_string(".red")?;
512            let red = ();
513            stream.expect_complete()?;
514            stream.expect_string(".sync")?;
515            let sync = ();
516            stream.expect_complete()?;
517            stream.expect_string(".aligned")?;
518            let aligned = ();
519            stream.expect_complete()?;
520            let shape4 = Shape4::parse(stream)?;
521            stream.expect_complete()?;
522            let num = Num::parse(stream)?;
523            stream.expect_complete()?;
524            let redop = Redop::parse(stream)?;
525            stream.expect_complete()?;
526            let saved_pos = stream.position();
527            let abs = stream.expect_string(".abs").is_ok();
528            if !abs {
529                stream.set_position(saved_pos);
530            }
531            stream.expect_complete()?;
532            let saved_pos = stream.position();
533            let nan = stream.expect_string(".NaN").is_ok();
534            if !nan {
535                stream.set_position(saved_pos);
536            }
537            stream.expect_complete()?;
538            stream.expect_string(".f32")?;
539            let f32 = ();
540            stream.expect_complete()?;
541            let r = GeneralOperand::parse(stream)?;
542            stream.expect_complete()?;
543            stream.expect(&PtxToken::Comma)?;
544            let redval = GeneralOperand::parse(stream)?;
545            stream.expect_complete()?;
546            stream.expect(&PtxToken::Comma)?;
547            let taddr = AddressOperand::parse(stream)?;
548            stream.expect_complete()?;
549            stream.expect(&PtxToken::Comma)?;
550            let immhalfsplitoff = GeneralOperand::parse(stream)?;
551            stream.expect_complete()?;
552            stream.expect_complete()?;
553            stream.expect(&PtxToken::Semicolon)?;
554            Ok(Tcgen05LdRedSyncAlignedShape4NumRedopAbsNanF32 {
555                ld,
556                red,
557                sync,
558                aligned,
559                shape4,
560                num,
561                redop,
562                abs,
563                nan,
564                f32,
565                r,
566                redval,
567                taddr,
568                immhalfsplitoff,
569            })
570        }
571    }
572
573    impl PtxParser for Tcgen05LdRedSyncAlignedShape3NumRedopType {
574        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
575            stream.expect_string("tcgen05")?;
576            stream.expect_string(".ld")?;
577            let ld = ();
578            stream.expect_complete()?;
579            stream.expect_string(".red")?;
580            let red = ();
581            stream.expect_complete()?;
582            stream.expect_string(".sync")?;
583            let sync = ();
584            stream.expect_complete()?;
585            stream.expect_string(".aligned")?;
586            let aligned = ();
587            stream.expect_complete()?;
588            let shape3 = Shape3::parse(stream)?;
589            stream.expect_complete()?;
590            let num = Num::parse(stream)?;
591            stream.expect_complete()?;
592            let redop = Redop::parse(stream)?;
593            stream.expect_complete()?;
594            let type_ = Type::parse(stream)?;
595            stream.expect_complete()?;
596            let r = GeneralOperand::parse(stream)?;
597            stream.expect_complete()?;
598            stream.expect(&PtxToken::Comma)?;
599            let redval = GeneralOperand::parse(stream)?;
600            stream.expect_complete()?;
601            stream.expect(&PtxToken::Comma)?;
602            let taddr = AddressOperand::parse(stream)?;
603            stream.expect_complete()?;
604            stream.expect_complete()?;
605            stream.expect(&PtxToken::Semicolon)?;
606            Ok(Tcgen05LdRedSyncAlignedShape3NumRedopType {
607                ld,
608                red,
609                sync,
610                aligned,
611                shape3,
612                num,
613                redop,
614                type_,
615                r,
616                redval,
617                taddr,
618            })
619        }
620    }
621
622    impl PtxParser for Tcgen05LdRedSyncAlignedShape4NumRedopType {
623        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
624            stream.expect_string("tcgen05")?;
625            stream.expect_string(".ld")?;
626            let ld = ();
627            stream.expect_complete()?;
628            stream.expect_string(".red")?;
629            let red = ();
630            stream.expect_complete()?;
631            stream.expect_string(".sync")?;
632            let sync = ();
633            stream.expect_complete()?;
634            stream.expect_string(".aligned")?;
635            let aligned = ();
636            stream.expect_complete()?;
637            let shape4 = Shape4::parse(stream)?;
638            stream.expect_complete()?;
639            let num = Num::parse(stream)?;
640            stream.expect_complete()?;
641            let redop = Redop::parse(stream)?;
642            stream.expect_complete()?;
643            let type_ = Type::parse(stream)?;
644            stream.expect_complete()?;
645            let r = GeneralOperand::parse(stream)?;
646            stream.expect_complete()?;
647            stream.expect(&PtxToken::Comma)?;
648            let redval = GeneralOperand::parse(stream)?;
649            stream.expect_complete()?;
650            stream.expect(&PtxToken::Comma)?;
651            let taddr = AddressOperand::parse(stream)?;
652            stream.expect_complete()?;
653            stream.expect(&PtxToken::Comma)?;
654            let immhalfsplitoff = GeneralOperand::parse(stream)?;
655            stream.expect_complete()?;
656            stream.expect_complete()?;
657            stream.expect(&PtxToken::Semicolon)?;
658            Ok(Tcgen05LdRedSyncAlignedShape4NumRedopType {
659                ld,
660                red,
661                sync,
662                aligned,
663                shape4,
664                num,
665                redop,
666                type_,
667                r,
668                redval,
669                taddr,
670                immhalfsplitoff,
671            })
672        }
673    }
674}
ptx_parser/parser/instruction/tcgen05_ld.rs

ptx_parser/parser/instruction/
tcgen05_ld.rs
