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::parser::{
24    PtxParseError, PtxParser, PtxTokenStream, Span,
25    util::{
26        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
27        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
28    },
29};
30use crate::r#type::common::*;
31use crate::{alt, ok, seq_n};
32
33pub mod section_0 {
34    use super::*;
35    use crate::r#type::instruction::tcgen05_ld::section_0::*;
36
37    // ============================================================================
38    // Generated enum parsers
39    // ============================================================================
40
41    impl PtxParser for Num {
42        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
43            alt!(
44                map(string_p(".x128"), |_, _span| Num::X128),
45                map(string_p(".x16"), |_, _span| Num::X16),
46                map(string_p(".x32"), |_, _span| Num::X32),
47                map(string_p(".x64"), |_, _span| Num::X64),
48                map(string_p(".x1"), |_, _span| Num::X1),
49                map(string_p(".x2"), |_, _span| Num::X2),
50                map(string_p(".x4"), |_, _span| Num::X4),
51                map(string_p(".x8"), |_, _span| Num::X8)
52            )
53        }
54    }
55
56    impl PtxParser for Pack {
57        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
58            alt!(map(string_p(".pack::16b"), |_, _span| Pack::Pack16b))
59        }
60    }
61
62    impl PtxParser for Redop {
63        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
64            alt!(
65                map(string_p(".min"), |_, _span| Redop::Min),
66                map(string_p(".max"), |_, _span| Redop::Max)
67            )
68        }
69    }
70
71    impl PtxParser for Shape1 {
72        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
73            alt!(
74                map(string_p(".16x128b"), |_, _span| Shape1::_16x128b),
75                map(string_p(".16x256b"), |_, _span| Shape1::_16x256b),
76                map(string_p(".16x64b"), |_, _span| Shape1::_16x64b),
77                map(string_p(".32x32b"), |_, _span| Shape1::_32x32b)
78            )
79        }
80    }
81
82    impl PtxParser for Shape2 {
83        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
84            alt!(map(string_p(".16x32bx2"), |_, _span| Shape2::_16x32bx2))
85        }
86    }
87
88    impl PtxParser for Shape3 {
89        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
90            alt!(map(string_p(".32x32b"), |_, _span| Shape3::_32x32b))
91        }
92    }
93
94    impl PtxParser for Shape4 {
95        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
96            alt!(map(string_p(".16x32bx2"), |_, _span| Shape4::_16x32bx2))
97        }
98    }
99
100    impl PtxParser for Type {
101        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
102            alt!(
103                map(string_p(".u32"), |_, _span| Type::U32),
104                map(string_p(".s32"), |_, _span| Type::S32)
105            )
106        }
107    }
108
109    impl PtxParser for Tcgen05LdSyncAlignedShape1NumPackB32 {
110        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
111            try_map(
112                seq_n!(
113                    string_p("tcgen05"),
114                    string_p(".ld"),
115                    string_p(".sync"),
116                    string_p(".aligned"),
117                    Shape1::parse(),
118                    Num::parse(),
119                    optional(Pack::parse()),
120                    string_p(".b32"),
121                    GeneralOperand::parse(),
122                    comma_p(),
123                    AddressOperand::parse(),
124                    semicolon_p()
125                ),
126                |(_, ld, sync, aligned, shape1, num, pack, b32, r, _, taddr, _), span| {
127                    ok!(Tcgen05LdSyncAlignedShape1NumPackB32 {
128                        ld = ld,
129                        sync = sync,
130                        aligned = aligned,
131                        shape1 = shape1,
132                        num = num,
133                        pack = pack,
134                        b32 = b32,
135                        r = r,
136                        taddr = taddr,
137
138                    })
139                },
140            )
141        }
142    }
143
144    impl PtxParser for Tcgen05LdSyncAlignedShape2NumPackB32 {
145        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
146            try_map(
147                seq_n!(
148                    string_p("tcgen05"),
149                    string_p(".ld"),
150                    string_p(".sync"),
151                    string_p(".aligned"),
152                    Shape2::parse(),
153                    Num::parse(),
154                    optional(Pack::parse()),
155                    string_p(".b32"),
156                    GeneralOperand::parse(),
157                    comma_p(),
158                    AddressOperand::parse(),
159                    comma_p(),
160                    GeneralOperand::parse(),
161                    semicolon_p()
162                ),
163                |(
164                    _,
165                    ld,
166                    sync,
167                    aligned,
168                    shape2,
169                    num,
170                    pack,
171                    b32,
172                    r,
173                    _,
174                    taddr,
175                    _,
176                    immhalfsplitoff,
177                    _,
178                ),
179                 span| {
180                    ok!(Tcgen05LdSyncAlignedShape2NumPackB32 {
181                        ld = ld,
182                        sync = sync,
183                        aligned = aligned,
184                        shape2 = shape2,
185                        num = num,
186                        pack = pack,
187                        b32 = b32,
188                        r = r,
189                        taddr = taddr,
190                        immhalfsplitoff = immhalfsplitoff,
191
192                    })
193                },
194            )
195        }
196    }
197
198    impl PtxParser for Tcgen05LdRedSyncAlignedShape3NumRedopAbsNanF32 {
199        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
200            try_map(
201                seq_n!(
202                    string_p("tcgen05"),
203                    string_p(".ld"),
204                    string_p(".red"),
205                    string_p(".sync"),
206                    string_p(".aligned"),
207                    Shape3::parse(),
208                    Num::parse(),
209                    Redop::parse(),
210                    map(optional(string_p(".abs")), |value, _| value.is_some()),
211                    map(optional(string_p(".NaN")), |value, _| value.is_some()),
212                    string_p(".f32"),
213                    GeneralOperand::parse(),
214                    comma_p(),
215                    GeneralOperand::parse(),
216                    comma_p(),
217                    AddressOperand::parse(),
218                    semicolon_p()
219                ),
220                |(
221                    _,
222                    ld,
223                    red,
224                    sync,
225                    aligned,
226                    shape3,
227                    num,
228                    redop,
229                    abs,
230                    nan,
231                    f32,
232                    r,
233                    _,
234                    redval,
235                    _,
236                    taddr,
237                    _,
238                ),
239                 span| {
240                    ok!(Tcgen05LdRedSyncAlignedShape3NumRedopAbsNanF32 {
241                        ld = ld,
242                        red = red,
243                        sync = sync,
244                        aligned = aligned,
245                        shape3 = shape3,
246                        num = num,
247                        redop = redop,
248                        abs = abs,
249                        nan = nan,
250                        f32 = f32,
251                        r = r,
252                        redval = redval,
253                        taddr = taddr,
254
255                    })
256                },
257            )
258        }
259    }
260
261    impl PtxParser for Tcgen05LdRedSyncAlignedShape4NumRedopAbsNanF32 {
262        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
263            try_map(
264                seq_n!(
265                    string_p("tcgen05"),
266                    string_p(".ld"),
267                    string_p(".red"),
268                    string_p(".sync"),
269                    string_p(".aligned"),
270                    Shape4::parse(),
271                    Num::parse(),
272                    Redop::parse(),
273                    map(optional(string_p(".abs")), |value, _| value.is_some()),
274                    map(optional(string_p(".NaN")), |value, _| value.is_some()),
275                    string_p(".f32"),
276                    GeneralOperand::parse(),
277                    comma_p(),
278                    GeneralOperand::parse(),
279                    comma_p(),
280                    AddressOperand::parse(),
281                    comma_p(),
282                    GeneralOperand::parse(),
283                    semicolon_p()
284                ),
285                |(
286                    _,
287                    ld,
288                    red,
289                    sync,
290                    aligned,
291                    shape4,
292                    num,
293                    redop,
294                    abs,
295                    nan,
296                    f32,
297                    r,
298                    _,
299                    redval,
300                    _,
301                    taddr,
302                    _,
303                    immhalfsplitoff,
304                    _,
305                ),
306                 span| {
307                    ok!(Tcgen05LdRedSyncAlignedShape4NumRedopAbsNanF32 {
308                        ld = ld,
309                        red = red,
310                        sync = sync,
311                        aligned = aligned,
312                        shape4 = shape4,
313                        num = num,
314                        redop = redop,
315                        abs = abs,
316                        nan = nan,
317                        f32 = f32,
318                        r = r,
319                        redval = redval,
320                        taddr = taddr,
321                        immhalfsplitoff = immhalfsplitoff,
322
323                    })
324                },
325            )
326        }
327    }
328
329    impl PtxParser for Tcgen05LdRedSyncAlignedShape3NumRedopType {
330        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
331            try_map(
332                seq_n!(
333                    string_p("tcgen05"),
334                    string_p(".ld"),
335                    string_p(".red"),
336                    string_p(".sync"),
337                    string_p(".aligned"),
338                    Shape3::parse(),
339                    Num::parse(),
340                    Redop::parse(),
341                    Type::parse(),
342                    GeneralOperand::parse(),
343                    comma_p(),
344                    GeneralOperand::parse(),
345                    comma_p(),
346                    AddressOperand::parse(),
347                    semicolon_p()
348                ),
349                |(
350                    _,
351                    ld,
352                    red,
353                    sync,
354                    aligned,
355                    shape3,
356                    num,
357                    redop,
358                    type_,
359                    r,
360                    _,
361                    redval,
362                    _,
363                    taddr,
364                    _,
365                ),
366                 span| {
367                    ok!(Tcgen05LdRedSyncAlignedShape3NumRedopType {
368                        ld = ld,
369                        red = red,
370                        sync = sync,
371                        aligned = aligned,
372                        shape3 = shape3,
373                        num = num,
374                        redop = redop,
375                        type_ = type_,
376                        r = r,
377                        redval = redval,
378                        taddr = taddr,
379
380                    })
381                },
382            )
383        }
384    }
385
386    impl PtxParser for Tcgen05LdRedSyncAlignedShape4NumRedopType {
387        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
388            try_map(
389                seq_n!(
390                    string_p("tcgen05"),
391                    string_p(".ld"),
392                    string_p(".red"),
393                    string_p(".sync"),
394                    string_p(".aligned"),
395                    Shape4::parse(),
396                    Num::parse(),
397                    Redop::parse(),
398                    Type::parse(),
399                    GeneralOperand::parse(),
400                    comma_p(),
401                    GeneralOperand::parse(),
402                    comma_p(),
403                    AddressOperand::parse(),
404                    comma_p(),
405                    GeneralOperand::parse(),
406                    semicolon_p()
407                ),
408                |(
409                    _,
410                    ld,
411                    red,
412                    sync,
413                    aligned,
414                    shape4,
415                    num,
416                    redop,
417                    type_,
418                    r,
419                    _,
420                    redval,
421                    _,
422                    taddr,
423                    _,
424                    immhalfsplitoff,
425                    _,
426                ),
427                 span| {
428                    ok!(Tcgen05LdRedSyncAlignedShape4NumRedopType {
429                        ld = ld,
430                        red = red,
431                        sync = sync,
432                        aligned = aligned,
433                        shape4 = shape4,
434                        num = num,
435                        redop = redop,
436                        type_ = type_,
437                        r = r,
438                        redval = redval,
439                        taddr = taddr,
440                        immhalfsplitoff = immhalfsplitoff,
441
442                    })
443                },
444            )
445        }
446    }
447}
ptx_parser/parser/instruction/tcgen05_ld.rs

ptx_parser/parser/instruction/
tcgen05_ld.rs
