ptx_parser/parser/instruction/
tcgen05_mma_sp.rs

1//! Original PTX specification:
2//!
3//! // 1. Floating-point type without block scaling:
4//! tcgen05.mma.sp.cta_group.kind  [d-tmem],  a-desc,  b-desc, [sp-meta-tmem] ,  idesc {, disable-output-lane }, enable-input-d{, scale-input-d};
5//! tcgen05.mma.sp.cta_group.kind  [d-tmem], [a-tmem], b-desc, [sp-meta-tmem] , idesc {, disable-output-lane }, enable-input-d{, scale-input-d};
6//! .kind       = { .kind::f16, .kind::tf32, .kind::f8f6f4 };
7//! .cta_group  = { .cta_group::1,  .cta_group::2 };
8//! ------------------------------------------------------------------
9//! // 2. Floating-point type with block scaling:
10//! tcgen05.mma.sp.cta_group.kind.block_scale{.scale_vectorsize} [d-tmem],  a-desc,  b-desc , [sp-meta-tmem] , idesc, [scale-A-tmem], [scale-B-tmem], enable-input-d;
11//! tcgen05.mma.sp.cta_group.kind.block_scale{.scale_vectorsize} [d-tmem], [a-tmem], b-desc , [sp-meta-tmem] , idesc, [scale-A-tmem], [scale-B-tmem], enable-input-d;
12//! .scale_vectorsize = { .scale_vec::1X, .scale_vec::2X, .scale_vec::4X, .block16, .block32 };
13//! .cta_group      = { .cta_group::1,  .cta_group::2 };
14//! .kind = { .kind::mxf8f6f4, .kind::mxf4, .kind::mxf4nvf4 };
15//! ------------------------------------------------------------------
16//! // 3. Convolution MMA with floating-point type without block scaling:
17//! tcgen05.mma.sp.cta_group.kind.collector_usage           [d-tmem],  a-desc,  b-desc, [sp-meta-tmem] ,  idesc {, disable-output-lane }, enable-input-d
18//! {, scale-input-d};
19//! tcgen05.mma.sp.cta_group.kind.ashift{.collector_usage}  [d-tmem], [a-tmem], b-desc, [sp-meta-tmem] , idesc {, disable-output-lane }, enable-input-d
20//! {, scale-input-d};
21//! tcgen05.mma.sp.cta_group.kind{.ashift}.collector_usage  [d-tmem], [a-tmem], b-desc, [sp-meta-tmem] , idesc {, disable-output-lane }, enable-input-d
22//! {, scale-input-d};
23//! .kind            = { .kind::f16, .kind::tf32, .kind::f8f6f4 };
24//! .collector_usage = { .collector::buffer::op };
25//! ::buffer         = { ::a };
26//! ::op             = { ::fill, ::use, ::lastuse, ::discard* };
27//! ------------------------------------------------------------------
28//! // 4. Activation Stationary MMA with floating-point type with block scaling:
29//! tcgen05.mma.sp.cta_group.kind.block_scale{.scale_vectorsize}.collector_usage [d-tmem],  a-desc,  b-desc , [sp-meta-tmem] , idesc, [scale-A-tmem], [scale-B-tmem], enable-input-d;
30//! tcgen05.mma.sp.cta_group.kind.block_scale{.scale_vectorsize}.collector_usage [d-tmem], [a-tmem], b-desc , [sp-meta-tmem] , idesc, [scale-A-tmem], [scale-B-tmem], enable-input-d;
31//! .kind = { .kind::mxf8f6f4, .kind::mxf4, .kind::mxf4nvf4 };
32//! .scale_vectorsize = { .scale_vec::1X, .scale_vec::2X, .scale_vec::4X, .block16, .block32 };
33//! .collector_usage = { .collector::buffer::op };
34//! ::buffer         = { ::a };
35//! ::op             = { ::fill, ::use, ::lastuse, ::discard* };
36//! ------------------------------------------------------------------
37//! // 5. Integer type:
38//! tcgen05.mma.sp.cta_group.kind::i8 [d-tmem],  a-desc,  b-desc, [sp-meta-tmem] , idesc {, disable-output-lane }, enable-input-d;
39//! tcgen05.mma.sp.cta_group.kind::i8 [d-tmem], [a-tmem], b-desc, [sp-meta-tmem] , idesc {, disable-output-lane }, enable-input-d;
40//! .cta_group      = { .cta_group::1,  .cta_group::2 };
41//! ------------------------------------------------------------------
42//! // 6. Convolution MMA with Integer type:
43//! tcgen05.mma.sp.cta_group.kind::i8.collector_usage          [d-tmem],  a-desc,  b-desc, [sp-meta-tmem], idesc {, disable-output-lane }, enable-input-d;
44//! tcgen05.mma.sp.cta_group.kind::i8.ashift{.collector_usage} [d-tmem], [a-tmem], b-desc, [sp-meta-tmem], idesc {, disable-output-lane }, enable-input-d;
45//! tcgen05.mma.sp.cta_group.kind::i8{.ashift}.collector_usage [d-tmem], [a-tmem], b-desc, [sp-meta-tmem], idesc {, disable-output-lane }, enable-input-d;
46//! .collector_usage = { .collector::buffer::op };
47//! ::buffer         = { ::a };
48//! ::op             = { ::fill, ::use, ::lastuse, ::discard* };
49
50#![allow(unused)]
51
52use crate::parser::{
53    PtxParseError, PtxParser, PtxTokenStream, Span,
54    util::{
55        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
56        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
57    },
58};
59use crate::r#type::common::*;
60use crate::{alt, ok, seq_n};
61
62pub mod section_0 {
63    use super::*;
64    use crate::r#type::instruction::tcgen05_mma_sp::section_0::*;
65
66    // ============================================================================
67    // Generated enum parsers
68    // ============================================================================
69
70    impl PtxParser for CtaGroup {
71        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
72            alt!(
73                map(string_p(".cta_group::1"), |_, _span| CtaGroup::CtaGroup1),
74                map(string_p(".cta_group::2"), |_, _span| CtaGroup::CtaGroup2)
75            )
76        }
77    }
78
79    impl PtxParser for Kind {
80        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
81            alt!(
82                map(string_p(".kind::f8f6f4"), |_, _span| Kind::KindF8f6f4),
83                map(string_p(".kind::tf32"), |_, _span| Kind::KindTf32),
84                map(string_p(".kind::f16"), |_, _span| Kind::KindF16)
85            )
86        }
87    }
88
89    impl PtxParser for Tcgen05MmaSpCtaGroupKind {
90        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
91            try_map(
92                seq_n!(
93                    string_p("tcgen05"),
94                    string_p(".mma"),
95                    string_p(".sp"),
96                    CtaGroup::parse(),
97                    Kind::parse(),
98                    AddressOperand::parse(),
99                    comma_p(),
100                    GeneralOperand::parse(),
101                    comma_p(),
102                    GeneralOperand::parse(),
103                    comma_p(),
104                    AddressOperand::parse(),
105                    comma_p(),
106                    GeneralOperand::parse(),
107                    map(
108                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
109                        |value, _| value.map(|(_, operand)| operand)
110                    ),
111                    comma_p(),
112                    GeneralOperand::parse(),
113                    map(
114                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
115                        |value, _| value.map(|(_, operand)| operand)
116                    ),
117                    semicolon_p()
118                ),
119                |(
120                    _,
121                    mma,
122                    sp,
123                    cta_group,
124                    kind,
125                    d_tmem,
126                    _,
127                    a_desc,
128                    _,
129                    b_desc,
130                    _,
131                    sp_meta_tmem,
132                    _,
133                    idesc,
134                    disable_output_lane,
135                    _,
136                    enable_input_d,
137                    scale_input_d,
138                    _,
139                ),
140                 span| {
141                    ok!(Tcgen05MmaSpCtaGroupKind {
142                        mma = mma,
143                        sp = sp,
144                        cta_group = cta_group,
145                        kind = kind,
146                        d_tmem = d_tmem,
147                        a_desc = a_desc,
148                        b_desc = b_desc,
149                        sp_meta_tmem = sp_meta_tmem,
150                        idesc = idesc,
151                        disable_output_lane = disable_output_lane,
152                        enable_input_d = enable_input_d,
153                        scale_input_d = scale_input_d,
154
155                    })
156                },
157            )
158        }
159    }
160
161    impl PtxParser for Tcgen05MmaSpCtaGroupKind1 {
162        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
163            try_map(
164                seq_n!(
165                    string_p("tcgen05"),
166                    string_p(".mma"),
167                    string_p(".sp"),
168                    CtaGroup::parse(),
169                    Kind::parse(),
170                    AddressOperand::parse(),
171                    comma_p(),
172                    AddressOperand::parse(),
173                    comma_p(),
174                    GeneralOperand::parse(),
175                    comma_p(),
176                    AddressOperand::parse(),
177                    comma_p(),
178                    GeneralOperand::parse(),
179                    map(
180                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
181                        |value, _| value.map(|(_, operand)| operand)
182                    ),
183                    comma_p(),
184                    GeneralOperand::parse(),
185                    map(
186                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
187                        |value, _| value.map(|(_, operand)| operand)
188                    ),
189                    semicolon_p()
190                ),
191                |(
192                    _,
193                    mma,
194                    sp,
195                    cta_group,
196                    kind,
197                    d_tmem,
198                    _,
199                    a_tmem,
200                    _,
201                    b_desc,
202                    _,
203                    sp_meta_tmem,
204                    _,
205                    idesc,
206                    disable_output_lane,
207                    _,
208                    enable_input_d,
209                    scale_input_d,
210                    _,
211                ),
212                 span| {
213                    ok!(Tcgen05MmaSpCtaGroupKind1 {
214                        mma = mma,
215                        sp = sp,
216                        cta_group = cta_group,
217                        kind = kind,
218                        d_tmem = d_tmem,
219                        a_tmem = a_tmem,
220                        b_desc = b_desc,
221                        sp_meta_tmem = sp_meta_tmem,
222                        idesc = idesc,
223                        disable_output_lane = disable_output_lane,
224                        enable_input_d = enable_input_d,
225                        scale_input_d = scale_input_d,
226
227                    })
228                },
229            )
230        }
231    }
232}
233
234pub mod section_1 {
235    use super::*;
236    use crate::r#type::instruction::tcgen05_mma_sp::section_1::*;
237
238    // ============================================================================
239    // Generated enum parsers
240    // ============================================================================
241
242    impl PtxParser for CtaGroup {
243        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
244            alt!(
245                map(string_p(".cta_group::1"), |_, _span| CtaGroup::CtaGroup1),
246                map(string_p(".cta_group::2"), |_, _span| CtaGroup::CtaGroup2)
247            )
248        }
249    }
250
251    impl PtxParser for Kind {
252        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
253            alt!(
254                map(string_p(".kind::mxf8f6f4"), |_, _span| Kind::KindMxf8f6f4),
255                map(string_p(".kind::mxf4nvf4"), |_, _span| Kind::KindMxf4nvf4),
256                map(string_p(".kind::mxf4"), |_, _span| Kind::KindMxf4)
257            )
258        }
259    }
260
261    impl PtxParser for ScaleVectorsize {
262        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
263            alt!(
264                map(string_p(".scale_vec::1X"), |_, _span| {
265                    ScaleVectorsize::ScaleVec1x
266                }),
267                map(string_p(".scale_vec::2X"), |_, _span| {
268                    ScaleVectorsize::ScaleVec2x
269                }),
270                map(string_p(".scale_vec::4X"), |_, _span| {
271                    ScaleVectorsize::ScaleVec4x
272                }),
273                map(string_p(".block16"), |_, _span| ScaleVectorsize::Block16),
274                map(string_p(".block32"), |_, _span| ScaleVectorsize::Block32)
275            )
276        }
277    }
278
279    impl PtxParser for Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsize {
280        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
281            try_map(
282                seq_n!(
283                    string_p("tcgen05"),
284                    string_p(".mma"),
285                    string_p(".sp"),
286                    CtaGroup::parse(),
287                    Kind::parse(),
288                    string_p(".block_scale"),
289                    optional(ScaleVectorsize::parse()),
290                    AddressOperand::parse(),
291                    comma_p(),
292                    GeneralOperand::parse(),
293                    comma_p(),
294                    GeneralOperand::parse(),
295                    comma_p(),
296                    AddressOperand::parse(),
297                    comma_p(),
298                    GeneralOperand::parse(),
299                    comma_p(),
300                    AddressOperand::parse(),
301                    comma_p(),
302                    AddressOperand::parse(),
303                    comma_p(),
304                    GeneralOperand::parse(),
305                    semicolon_p()
306                ),
307                |(
308                    _,
309                    mma,
310                    sp,
311                    cta_group,
312                    kind,
313                    block_scale,
314                    scale_vectorsize,
315                    d_tmem,
316                    _,
317                    a_desc,
318                    _,
319                    b_desc,
320                    _,
321                    sp_meta_tmem,
322                    _,
323                    idesc,
324                    _,
325                    scale_a_tmem,
326                    _,
327                    scale_b_tmem,
328                    _,
329                    enable_input_d,
330                    _,
331                ),
332                 span| {
333                    ok!(Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsize {
334                        mma = mma,
335                        sp = sp,
336                        cta_group = cta_group,
337                        kind = kind,
338                        block_scale = block_scale,
339                        scale_vectorsize = scale_vectorsize,
340                        d_tmem = d_tmem,
341                        a_desc = a_desc,
342                        b_desc = b_desc,
343                        sp_meta_tmem = sp_meta_tmem,
344                        idesc = idesc,
345                        scale_a_tmem = scale_a_tmem,
346                        scale_b_tmem = scale_b_tmem,
347                        enable_input_d = enable_input_d,
348
349                    })
350                },
351            )
352        }
353    }
354
355    impl PtxParser for Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsize1 {
356        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
357            try_map(
358                seq_n!(
359                    string_p("tcgen05"),
360                    string_p(".mma"),
361                    string_p(".sp"),
362                    CtaGroup::parse(),
363                    Kind::parse(),
364                    string_p(".block_scale"),
365                    optional(ScaleVectorsize::parse()),
366                    AddressOperand::parse(),
367                    comma_p(),
368                    AddressOperand::parse(),
369                    comma_p(),
370                    GeneralOperand::parse(),
371                    comma_p(),
372                    AddressOperand::parse(),
373                    comma_p(),
374                    GeneralOperand::parse(),
375                    comma_p(),
376                    AddressOperand::parse(),
377                    comma_p(),
378                    AddressOperand::parse(),
379                    comma_p(),
380                    GeneralOperand::parse(),
381                    semicolon_p()
382                ),
383                |(
384                    _,
385                    mma,
386                    sp,
387                    cta_group,
388                    kind,
389                    block_scale,
390                    scale_vectorsize,
391                    d_tmem,
392                    _,
393                    a_tmem,
394                    _,
395                    b_desc,
396                    _,
397                    sp_meta_tmem,
398                    _,
399                    idesc,
400                    _,
401                    scale_a_tmem,
402                    _,
403                    scale_b_tmem,
404                    _,
405                    enable_input_d,
406                    _,
407                ),
408                 span| {
409                    ok!(Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsize1 {
410                        mma = mma,
411                        sp = sp,
412                        cta_group = cta_group,
413                        kind = kind,
414                        block_scale = block_scale,
415                        scale_vectorsize = scale_vectorsize,
416                        d_tmem = d_tmem,
417                        a_tmem = a_tmem,
418                        b_desc = b_desc,
419                        sp_meta_tmem = sp_meta_tmem,
420                        idesc = idesc,
421                        scale_a_tmem = scale_a_tmem,
422                        scale_b_tmem = scale_b_tmem,
423                        enable_input_d = enable_input_d,
424
425                    })
426                },
427            )
428        }
429    }
430}
431
432pub mod section_2 {
433    use super::*;
434    use crate::r#type::instruction::tcgen05_mma_sp::section_2::*;
435
436    // ============================================================================
437    // Generated enum parsers
438    // ============================================================================
439
440    impl PtxParser for CollectorUsage {
441        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
442            alt!(map(
443                |stream| {
444                    stream.try_with_span(|stream| {
445                        stream.with_partial_token_mode(|stream| {
446                            stream.expect_string(".collector")?;
447                            let part0 = match stream.expect_strings(&["::a"])? {
448                                0 => Buffer::A,
449                                _ => unreachable!(),
450                            };
451                            let part1 = match stream.expect_strings(&[
452                                "::discard*",
453                                "::lastuse",
454                                "::fill",
455                                "::use",
456                            ])? {
457                                0 => Op::Discard,
458                                1 => Op::Lastuse,
459                                2 => Op::Fill,
460                                3 => Op::Use,
461                                _ => unreachable!(),
462                            };
463                            Ok(((), part0, part1))
464                        })
465                    })
466                },
467                |(collector, buffer, op), _span| CollectorUsage::CollectorBufferOp(
468                    collector, buffer, op
469                )
470            ))
471        }
472    }
473
474    impl PtxParser for CtaGroup {
475        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
476            alt!(
477                map(string_p(".cta_group::1"), |_, _span| CtaGroup::CtaGroup1),
478                map(string_p(".cta_group::2"), |_, _span| CtaGroup::CtaGroup2)
479            )
480        }
481    }
482
483    impl PtxParser for Kind {
484        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
485            alt!(
486                map(string_p(".kind::f8f6f4"), |_, _span| Kind::KindF8f6f4),
487                map(string_p(".kind::tf32"), |_, _span| Kind::KindTf32),
488                map(string_p(".kind::f16"), |_, _span| Kind::KindF16)
489            )
490        }
491    }
492
493    impl PtxParser for Tcgen05MmaSpCtaGroupKindCollectorUsage {
494        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
495            try_map(
496                seq_n!(
497                    string_p("tcgen05"),
498                    string_p(".mma"),
499                    string_p(".sp"),
500                    CtaGroup::parse(),
501                    Kind::parse(),
502                    CollectorUsage::parse(),
503                    AddressOperand::parse(),
504                    comma_p(),
505                    GeneralOperand::parse(),
506                    comma_p(),
507                    GeneralOperand::parse(),
508                    comma_p(),
509                    AddressOperand::parse(),
510                    comma_p(),
511                    GeneralOperand::parse(),
512                    map(
513                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
514                        |value, _| value.map(|(_, operand)| operand)
515                    ),
516                    comma_p(),
517                    GeneralOperand::parse(),
518                    map(
519                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
520                        |value, _| value.map(|(_, operand)| operand)
521                    ),
522                    semicolon_p()
523                ),
524                |(
525                    _,
526                    mma,
527                    sp,
528                    cta_group,
529                    kind,
530                    collector_usage,
531                    d_tmem,
532                    _,
533                    a_desc,
534                    _,
535                    b_desc,
536                    _,
537                    sp_meta_tmem,
538                    _,
539                    idesc,
540                    disable_output_lane,
541                    _,
542                    enable_input_d,
543                    scale_input_d,
544                    _,
545                ),
546                 span| {
547                    ok!(Tcgen05MmaSpCtaGroupKindCollectorUsage {
548                        mma = mma,
549                        sp = sp,
550                        cta_group = cta_group,
551                        kind = kind,
552                        collector_usage = collector_usage,
553                        d_tmem = d_tmem,
554                        a_desc = a_desc,
555                        b_desc = b_desc,
556                        sp_meta_tmem = sp_meta_tmem,
557                        idesc = idesc,
558                        disable_output_lane = disable_output_lane,
559                        enable_input_d = enable_input_d,
560                        scale_input_d = scale_input_d,
561
562                    })
563                },
564            )
565        }
566    }
567
568    impl PtxParser for Tcgen05MmaSpCtaGroupKindAshiftCollectorUsage {
569        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
570            try_map(
571                seq_n!(
572                    string_p("tcgen05"),
573                    string_p(".mma"),
574                    string_p(".sp"),
575                    CtaGroup::parse(),
576                    Kind::parse(),
577                    string_p(".ashift"),
578                    optional(CollectorUsage::parse()),
579                    AddressOperand::parse(),
580                    comma_p(),
581                    AddressOperand::parse(),
582                    comma_p(),
583                    GeneralOperand::parse(),
584                    comma_p(),
585                    AddressOperand::parse(),
586                    comma_p(),
587                    GeneralOperand::parse(),
588                    map(
589                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
590                        |value, _| value.map(|(_, operand)| operand)
591                    ),
592                    comma_p(),
593                    GeneralOperand::parse(),
594                    map(
595                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
596                        |value, _| value.map(|(_, operand)| operand)
597                    ),
598                    semicolon_p()
599                ),
600                |(
601                    _,
602                    mma,
603                    sp,
604                    cta_group,
605                    kind,
606                    ashift,
607                    collector_usage,
608                    d_tmem,
609                    _,
610                    a_tmem,
611                    _,
612                    b_desc,
613                    _,
614                    sp_meta_tmem,
615                    _,
616                    idesc,
617                    disable_output_lane,
618                    _,
619                    enable_input_d,
620                    scale_input_d,
621                    _,
622                ),
623                 span| {
624                    ok!(Tcgen05MmaSpCtaGroupKindAshiftCollectorUsage {
625                        mma = mma,
626                        sp = sp,
627                        cta_group = cta_group,
628                        kind = kind,
629                        ashift = ashift,
630                        collector_usage = collector_usage,
631                        d_tmem = d_tmem,
632                        a_tmem = a_tmem,
633                        b_desc = b_desc,
634                        sp_meta_tmem = sp_meta_tmem,
635                        idesc = idesc,
636                        disable_output_lane = disable_output_lane,
637                        enable_input_d = enable_input_d,
638                        scale_input_d = scale_input_d,
639
640                    })
641                },
642            )
643        }
644    }
645
646    impl PtxParser for Tcgen05MmaSpCtaGroupKindAshiftCollectorUsage1 {
647        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
648            try_map(
649                seq_n!(
650                    string_p("tcgen05"),
651                    string_p(".mma"),
652                    string_p(".sp"),
653                    CtaGroup::parse(),
654                    Kind::parse(),
655                    map(optional(string_p(".ashift")), |value, _| value.is_some()),
656                    CollectorUsage::parse(),
657                    AddressOperand::parse(),
658                    comma_p(),
659                    AddressOperand::parse(),
660                    comma_p(),
661                    GeneralOperand::parse(),
662                    comma_p(),
663                    AddressOperand::parse(),
664                    comma_p(),
665                    GeneralOperand::parse(),
666                    map(
667                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
668                        |value, _| value.map(|(_, operand)| operand)
669                    ),
670                    comma_p(),
671                    GeneralOperand::parse(),
672                    map(
673                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
674                        |value, _| value.map(|(_, operand)| operand)
675                    ),
676                    semicolon_p()
677                ),
678                |(
679                    _,
680                    mma,
681                    sp,
682                    cta_group,
683                    kind,
684                    ashift,
685                    collector_usage,
686                    d_tmem,
687                    _,
688                    a_tmem,
689                    _,
690                    b_desc,
691                    _,
692                    sp_meta_tmem,
693                    _,
694                    idesc,
695                    disable_output_lane,
696                    _,
697                    enable_input_d,
698                    scale_input_d,
699                    _,
700                ),
701                 span| {
702                    ok!(Tcgen05MmaSpCtaGroupKindAshiftCollectorUsage1 {
703                        mma = mma,
704                        sp = sp,
705                        cta_group = cta_group,
706                        kind = kind,
707                        ashift = ashift,
708                        collector_usage = collector_usage,
709                        d_tmem = d_tmem,
710                        a_tmem = a_tmem,
711                        b_desc = b_desc,
712                        sp_meta_tmem = sp_meta_tmem,
713                        idesc = idesc,
714                        disable_output_lane = disable_output_lane,
715                        enable_input_d = enable_input_d,
716                        scale_input_d = scale_input_d,
717
718                    })
719                },
720            )
721        }
722    }
723}
724
725pub mod section_3 {
726    use super::*;
727    use crate::r#type::instruction::tcgen05_mma_sp::section_3::*;
728
729    // ============================================================================
730    // Generated enum parsers
731    // ============================================================================
732
733    impl PtxParser for CollectorUsage {
734        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
735            alt!(map(
736                |stream| {
737                    stream.try_with_span(|stream| {
738                        stream.with_partial_token_mode(|stream| {
739                            stream.expect_string(".collector")?;
740                            let part0 = match stream.expect_strings(&["::a"])? {
741                                0 => Buffer::A,
742                                _ => unreachable!(),
743                            };
744                            let part1 = match stream.expect_strings(&[
745                                "::discard*",
746                                "::lastuse",
747                                "::fill",
748                                "::use",
749                            ])? {
750                                0 => Op::Discard,
751                                1 => Op::Lastuse,
752                                2 => Op::Fill,
753                                3 => Op::Use,
754                                _ => unreachable!(),
755                            };
756                            Ok(((), part0, part1))
757                        })
758                    })
759                },
760                |(collector, buffer, op), _span| CollectorUsage::CollectorBufferOp(
761                    collector, buffer, op
762                )
763            ))
764        }
765    }
766
767    impl PtxParser for CtaGroup {
768        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
769            alt!(
770                map(string_p(".cta_group::1"), |_, _span| CtaGroup::CtaGroup1),
771                map(string_p(".cta_group::2"), |_, _span| CtaGroup::CtaGroup2)
772            )
773        }
774    }
775
776    impl PtxParser for Kind {
777        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
778            alt!(
779                map(string_p(".kind::mxf8f6f4"), |_, _span| Kind::KindMxf8f6f4),
780                map(string_p(".kind::mxf4nvf4"), |_, _span| Kind::KindMxf4nvf4),
781                map(string_p(".kind::mxf4"), |_, _span| Kind::KindMxf4)
782            )
783        }
784    }
785
786    impl PtxParser for ScaleVectorsize {
787        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
788            alt!(
789                map(string_p(".scale_vec::1X"), |_, _span| {
790                    ScaleVectorsize::ScaleVec1x
791                }),
792                map(string_p(".scale_vec::2X"), |_, _span| {
793                    ScaleVectorsize::ScaleVec2x
794                }),
795                map(string_p(".scale_vec::4X"), |_, _span| {
796                    ScaleVectorsize::ScaleVec4x
797                }),
798                map(string_p(".block16"), |_, _span| ScaleVectorsize::Block16),
799                map(string_p(".block32"), |_, _span| ScaleVectorsize::Block32)
800            )
801        }
802    }
803
804    impl PtxParser for Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage {
805        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
806            try_map(
807                seq_n!(
808                    string_p("tcgen05"),
809                    string_p(".mma"),
810                    string_p(".sp"),
811                    CtaGroup::parse(),
812                    Kind::parse(),
813                    string_p(".block_scale"),
814                    optional(ScaleVectorsize::parse()),
815                    CollectorUsage::parse(),
816                    AddressOperand::parse(),
817                    comma_p(),
818                    GeneralOperand::parse(),
819                    comma_p(),
820                    GeneralOperand::parse(),
821                    comma_p(),
822                    AddressOperand::parse(),
823                    comma_p(),
824                    GeneralOperand::parse(),
825                    comma_p(),
826                    AddressOperand::parse(),
827                    comma_p(),
828                    AddressOperand::parse(),
829                    comma_p(),
830                    GeneralOperand::parse(),
831                    semicolon_p()
832                ),
833                |(
834                    _,
835                    mma,
836                    sp,
837                    cta_group,
838                    kind,
839                    block_scale,
840                    scale_vectorsize,
841                    collector_usage,
842                    d_tmem,
843                    _,
844                    a_desc,
845                    _,
846                    b_desc,
847                    _,
848                    sp_meta_tmem,
849                    _,
850                    idesc,
851                    _,
852                    scale_a_tmem,
853                    _,
854                    scale_b_tmem,
855                    _,
856                    enable_input_d,
857                    _,
858                ),
859                 span| {
860                    ok!(Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage {
861                        mma = mma,
862                        sp = sp,
863                        cta_group = cta_group,
864                        kind = kind,
865                        block_scale = block_scale,
866                        scale_vectorsize = scale_vectorsize,
867                        collector_usage = collector_usage,
868                        d_tmem = d_tmem,
869                        a_desc = a_desc,
870                        b_desc = b_desc,
871                        sp_meta_tmem = sp_meta_tmem,
872                        idesc = idesc,
873                        scale_a_tmem = scale_a_tmem,
874                        scale_b_tmem = scale_b_tmem,
875                        enable_input_d = enable_input_d,
876
877                    })
878                },
879            )
880        }
881    }
882
883    impl PtxParser for Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage1 {
884        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
885            try_map(
886                seq_n!(
887                    string_p("tcgen05"),
888                    string_p(".mma"),
889                    string_p(".sp"),
890                    CtaGroup::parse(),
891                    Kind::parse(),
892                    string_p(".block_scale"),
893                    optional(ScaleVectorsize::parse()),
894                    CollectorUsage::parse(),
895                    AddressOperand::parse(),
896                    comma_p(),
897                    AddressOperand::parse(),
898                    comma_p(),
899                    GeneralOperand::parse(),
900                    comma_p(),
901                    AddressOperand::parse(),
902                    comma_p(),
903                    GeneralOperand::parse(),
904                    comma_p(),
905                    AddressOperand::parse(),
906                    comma_p(),
907                    AddressOperand::parse(),
908                    comma_p(),
909                    GeneralOperand::parse(),
910                    semicolon_p()
911                ),
912                |(
913                    _,
914                    mma,
915                    sp,
916                    cta_group,
917                    kind,
918                    block_scale,
919                    scale_vectorsize,
920                    collector_usage,
921                    d_tmem,
922                    _,
923                    a_tmem,
924                    _,
925                    b_desc,
926                    _,
927                    sp_meta_tmem,
928                    _,
929                    idesc,
930                    _,
931                    scale_a_tmem,
932                    _,
933                    scale_b_tmem,
934                    _,
935                    enable_input_d,
936                    _,
937                ),
938                 span| {
939                    ok!(Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage1 {
940                        mma = mma,
941                        sp = sp,
942                        cta_group = cta_group,
943                        kind = kind,
944                        block_scale = block_scale,
945                        scale_vectorsize = scale_vectorsize,
946                        collector_usage = collector_usage,
947                        d_tmem = d_tmem,
948                        a_tmem = a_tmem,
949                        b_desc = b_desc,
950                        sp_meta_tmem = sp_meta_tmem,
951                        idesc = idesc,
952                        scale_a_tmem = scale_a_tmem,
953                        scale_b_tmem = scale_b_tmem,
954                        enable_input_d = enable_input_d,
955
956                    })
957                },
958            )
959        }
960    }
961}
962
963pub mod section_4 {
964    use super::*;
965    use crate::r#type::instruction::tcgen05_mma_sp::section_4::*;
966
967    // ============================================================================
968    // Generated enum parsers
969    // ============================================================================
970
971    impl PtxParser for CtaGroup {
972        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
973            alt!(
974                map(string_p(".cta_group::1"), |_, _span| CtaGroup::CtaGroup1),
975                map(string_p(".cta_group::2"), |_, _span| CtaGroup::CtaGroup2)
976            )
977        }
978    }
979
980    impl PtxParser for Tcgen05MmaSpCtaGroupKindI8 {
981        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
982            try_map(
983                seq_n!(
984                    string_p("tcgen05"),
985                    string_p(".mma"),
986                    string_p(".sp"),
987                    CtaGroup::parse(),
988                    string_p(".kind::i8"),
989                    AddressOperand::parse(),
990                    comma_p(),
991                    GeneralOperand::parse(),
992                    comma_p(),
993                    GeneralOperand::parse(),
994                    comma_p(),
995                    AddressOperand::parse(),
996                    comma_p(),
997                    GeneralOperand::parse(),
998                    map(
999                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
1000                        |value, _| value.map(|(_, operand)| operand)
1001                    ),
1002                    comma_p(),
1003                    GeneralOperand::parse(),
1004                    semicolon_p()
1005                ),
1006                |(
1007                    _,
1008                    mma,
1009                    sp,
1010                    cta_group,
1011                    kind_i8,
1012                    d_tmem,
1013                    _,
1014                    a_desc,
1015                    _,
1016                    b_desc,
1017                    _,
1018                    sp_meta_tmem,
1019                    _,
1020                    idesc,
1021                    disable_output_lane,
1022                    _,
1023                    enable_input_d,
1024                    _,
1025                ),
1026                 span| {
1027                    ok!(Tcgen05MmaSpCtaGroupKindI8 {
1028                        mma = mma,
1029                        sp = sp,
1030                        cta_group = cta_group,
1031                        kind_i8 = kind_i8,
1032                        d_tmem = d_tmem,
1033                        a_desc = a_desc,
1034                        b_desc = b_desc,
1035                        sp_meta_tmem = sp_meta_tmem,
1036                        idesc = idesc,
1037                        disable_output_lane = disable_output_lane,
1038                        enable_input_d = enable_input_d,
1039
1040                    })
1041                },
1042            )
1043        }
1044    }
1045
1046    impl PtxParser for Tcgen05MmaSpCtaGroupKindI81 {
1047        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
1048            try_map(
1049                seq_n!(
1050                    string_p("tcgen05"),
1051                    string_p(".mma"),
1052                    string_p(".sp"),
1053                    CtaGroup::parse(),
1054                    string_p(".kind::i8"),
1055                    AddressOperand::parse(),
1056                    comma_p(),
1057                    AddressOperand::parse(),
1058                    comma_p(),
1059                    GeneralOperand::parse(),
1060                    comma_p(),
1061                    AddressOperand::parse(),
1062                    comma_p(),
1063                    GeneralOperand::parse(),
1064                    map(
1065                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
1066                        |value, _| value.map(|(_, operand)| operand)
1067                    ),
1068                    comma_p(),
1069                    GeneralOperand::parse(),
1070                    semicolon_p()
1071                ),
1072                |(
1073                    _,
1074                    mma,
1075                    sp,
1076                    cta_group,
1077                    kind_i8,
1078                    d_tmem,
1079                    _,
1080                    a_tmem,
1081                    _,
1082                    b_desc,
1083                    _,
1084                    sp_meta_tmem,
1085                    _,
1086                    idesc,
1087                    disable_output_lane,
1088                    _,
1089                    enable_input_d,
1090                    _,
1091                ),
1092                 span| {
1093                    ok!(Tcgen05MmaSpCtaGroupKindI81 {
1094                        mma = mma,
1095                        sp = sp,
1096                        cta_group = cta_group,
1097                        kind_i8 = kind_i8,
1098                        d_tmem = d_tmem,
1099                        a_tmem = a_tmem,
1100                        b_desc = b_desc,
1101                        sp_meta_tmem = sp_meta_tmem,
1102                        idesc = idesc,
1103                        disable_output_lane = disable_output_lane,
1104                        enable_input_d = enable_input_d,
1105
1106                    })
1107                },
1108            )
1109        }
1110    }
1111}
1112
1113pub mod section_5 {
1114    use super::*;
1115    use crate::r#type::instruction::tcgen05_mma_sp::section_5::*;
1116
1117    // ============================================================================
1118    // Generated enum parsers
1119    // ============================================================================
1120
1121    impl PtxParser for CollectorUsage {
1122        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
1123            alt!(map(
1124                |stream| {
1125                    stream.try_with_span(|stream| {
1126                        stream.with_partial_token_mode(|stream| {
1127                            stream.expect_string(".collector")?;
1128                            let part0 = match stream.expect_strings(&["::a"])? {
1129                                0 => Buffer::A,
1130                                _ => unreachable!(),
1131                            };
1132                            let part1 = match stream.expect_strings(&[
1133                                "::discard*",
1134                                "::lastuse",
1135                                "::fill",
1136                                "::use",
1137                            ])? {
1138                                0 => Op::Discard,
1139                                1 => Op::Lastuse,
1140                                2 => Op::Fill,
1141                                3 => Op::Use,
1142                                _ => unreachable!(),
1143                            };
1144                            Ok(((), part0, part1))
1145                        })
1146                    })
1147                },
1148                |(collector, buffer, op), _span| CollectorUsage::CollectorBufferOp(
1149                    collector, buffer, op
1150                )
1151            ))
1152        }
1153    }
1154
1155    impl PtxParser for CtaGroup {
1156        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
1157            alt!(
1158                map(string_p(".cta_group::1"), |_, _span| CtaGroup::CtaGroup1),
1159                map(string_p(".cta_group::2"), |_, _span| CtaGroup::CtaGroup2)
1160            )
1161        }
1162    }
1163
1164    impl PtxParser for Tcgen05MmaSpCtaGroupKindI8CollectorUsage {
1165        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
1166            try_map(
1167                seq_n!(
1168                    string_p("tcgen05"),
1169                    string_p(".mma"),
1170                    string_p(".sp"),
1171                    CtaGroup::parse(),
1172                    string_p(".kind::i8"),
1173                    CollectorUsage::parse(),
1174                    AddressOperand::parse(),
1175                    comma_p(),
1176                    GeneralOperand::parse(),
1177                    comma_p(),
1178                    GeneralOperand::parse(),
1179                    comma_p(),
1180                    AddressOperand::parse(),
1181                    comma_p(),
1182                    GeneralOperand::parse(),
1183                    map(
1184                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
1185                        |value, _| value.map(|(_, operand)| operand)
1186                    ),
1187                    comma_p(),
1188                    GeneralOperand::parse(),
1189                    semicolon_p()
1190                ),
1191                |(
1192                    _,
1193                    mma,
1194                    sp,
1195                    cta_group,
1196                    kind_i8,
1197                    collector_usage,
1198                    d_tmem,
1199                    _,
1200                    a_desc,
1201                    _,
1202                    b_desc,
1203                    _,
1204                    sp_meta_tmem,
1205                    _,
1206                    idesc,
1207                    disable_output_lane,
1208                    _,
1209                    enable_input_d,
1210                    _,
1211                ),
1212                 span| {
1213                    ok!(Tcgen05MmaSpCtaGroupKindI8CollectorUsage {
1214                        mma = mma,
1215                        sp = sp,
1216                        cta_group = cta_group,
1217                        kind_i8 = kind_i8,
1218                        collector_usage = collector_usage,
1219                        d_tmem = d_tmem,
1220                        a_desc = a_desc,
1221                        b_desc = b_desc,
1222                        sp_meta_tmem = sp_meta_tmem,
1223                        idesc = idesc,
1224                        disable_output_lane = disable_output_lane,
1225                        enable_input_d = enable_input_d,
1226
1227                    })
1228                },
1229            )
1230        }
1231    }
1232
1233    impl PtxParser for Tcgen05MmaSpCtaGroupKindI8AshiftCollectorUsage {
1234        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
1235            try_map(
1236                seq_n!(
1237                    string_p("tcgen05"),
1238                    string_p(".mma"),
1239                    string_p(".sp"),
1240                    CtaGroup::parse(),
1241                    string_p(".kind::i8"),
1242                    string_p(".ashift"),
1243                    optional(CollectorUsage::parse()),
1244                    AddressOperand::parse(),
1245                    comma_p(),
1246                    AddressOperand::parse(),
1247                    comma_p(),
1248                    GeneralOperand::parse(),
1249                    comma_p(),
1250                    AddressOperand::parse(),
1251                    comma_p(),
1252                    GeneralOperand::parse(),
1253                    map(
1254                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
1255                        |value, _| value.map(|(_, operand)| operand)
1256                    ),
1257                    comma_p(),
1258                    GeneralOperand::parse(),
1259                    semicolon_p()
1260                ),
1261                |(
1262                    _,
1263                    mma,
1264                    sp,
1265                    cta_group,
1266                    kind_i8,
1267                    ashift,
1268                    collector_usage,
1269                    d_tmem,
1270                    _,
1271                    a_tmem,
1272                    _,
1273                    b_desc,
1274                    _,
1275                    sp_meta_tmem,
1276                    _,
1277                    idesc,
1278                    disable_output_lane,
1279                    _,
1280                    enable_input_d,
1281                    _,
1282                ),
1283                 span| {
1284                    ok!(Tcgen05MmaSpCtaGroupKindI8AshiftCollectorUsage {
1285                        mma = mma,
1286                        sp = sp,
1287                        cta_group = cta_group,
1288                        kind_i8 = kind_i8,
1289                        ashift = ashift,
1290                        collector_usage = collector_usage,
1291                        d_tmem = d_tmem,
1292                        a_tmem = a_tmem,
1293                        b_desc = b_desc,
1294                        sp_meta_tmem = sp_meta_tmem,
1295                        idesc = idesc,
1296                        disable_output_lane = disable_output_lane,
1297                        enable_input_d = enable_input_d,
1298
1299                    })
1300                },
1301            )
1302        }
1303    }
1304
1305    impl PtxParser for Tcgen05MmaSpCtaGroupKindI8AshiftCollectorUsage1 {
1306        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
1307            try_map(
1308                seq_n!(
1309                    string_p("tcgen05"),
1310                    string_p(".mma"),
1311                    string_p(".sp"),
1312                    CtaGroup::parse(),
1313                    string_p(".kind::i8"),
1314                    map(optional(string_p(".ashift")), |value, _| value.is_some()),
1315                    CollectorUsage::parse(),
1316                    AddressOperand::parse(),
1317                    comma_p(),
1318                    AddressOperand::parse(),
1319                    comma_p(),
1320                    GeneralOperand::parse(),
1321                    comma_p(),
1322                    AddressOperand::parse(),
1323                    comma_p(),
1324                    GeneralOperand::parse(),
1325                    map(
1326                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
1327                        |value, _| value.map(|(_, operand)| operand)
1328                    ),
1329                    comma_p(),
1330                    GeneralOperand::parse(),
1331                    semicolon_p()
1332                ),
1333                |(
1334                    _,
1335                    mma,
1336                    sp,
1337                    cta_group,
1338                    kind_i8,
1339                    ashift,
1340                    collector_usage,
1341                    d_tmem,
1342                    _,
1343                    a_tmem,
1344                    _,
1345                    b_desc,
1346                    _,
1347                    sp_meta_tmem,
1348                    _,
1349                    idesc,
1350                    disable_output_lane,
1351                    _,
1352                    enable_input_d,
1353                    _,
1354                ),
1355                 span| {
1356                    ok!(Tcgen05MmaSpCtaGroupKindI8AshiftCollectorUsage1 {
1357                        mma = mma,
1358                        sp = sp,
1359                        cta_group = cta_group,
1360                        kind_i8 = kind_i8,
1361                        ashift = ashift,
1362                        collector_usage = collector_usage,
1363                        d_tmem = d_tmem,
1364                        a_tmem = a_tmem,
1365                        b_desc = b_desc,
1366                        sp_meta_tmem = sp_meta_tmem,
1367                        idesc = idesc,
1368                        disable_output_lane = disable_output_lane,
1369                        enable_input_d = enable_input_d,
1370
1371                    })
1372                },
1373            )
1374        }
1375    }
1376}
ptx_parser/parser/instruction/tcgen05_mma_sp.rs

ptx_parser/parser/instruction/
tcgen05_mma_sp.rs
