ptx_parser/parser/instruction/
tcgen05_cp.rs

1//! Original PTX specification:
2//!
3//! tcgen05.cp.cta_group.shape{.multicast}{.dst_src_fmt} [taddr], s-desc;
4//! .cta_group = { .cta_group::1, .cta_group::2 };
5//! .dst_src_fmt   = { .b8x16.b6x16_p32, .b8x16.b4x16_p64 };
6//! .shape     = { .128x256b, .4x256b, .128x128b, .64x128b**, .32x128b*** };
7//! .multicast = { .warpx2::02_13** , .warpx2::01_23**, .warpx4*** };
8
9#![allow(unused)]
10
11use crate::parser::{
12    PtxParseError, PtxParser, PtxTokenStream, Span,
13    util::{
14        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
15        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
16    },
17};
18use crate::r#type::common::*;
19use crate::{alt, ok, seq_n};
20
21pub mod section_0 {
22    use super::*;
23    use crate::r#type::instruction::tcgen05_cp::section_0::*;
24
25    // ============================================================================
26    // Generated enum parsers
27    // ============================================================================
28
29    impl PtxParser for CtaGroup {
30        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
31            alt!(
32                map(string_p(".cta_group::1"), |_, _span| CtaGroup::CtaGroup1),
33                map(string_p(".cta_group::2"), |_, _span| CtaGroup::CtaGroup2)
34            )
35        }
36    }
37
38    impl PtxParser for DstSrcFmt {
39        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
40            alt!(
41                map(string_p(".b8x16.b6x16_p32"), |_, _span| {
42                    DstSrcFmt::B8x16B6x16P32
43                }),
44                map(string_p(".b8x16.b4x16_p64"), |_, _span| {
45                    DstSrcFmt::B8x16B4x16P64
46                })
47            )
48        }
49    }
50
51    impl PtxParser for Multicast {
52        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
53            alt!(
54                map(string_p(".warpx2::02_13**"), |_, _span| {
55                    Multicast::Warpx20213
56                }),
57                map(string_p(".warpx2::01_23**"), |_, _span| {
58                    Multicast::Warpx20123
59                }),
60                map(string_p(".warpx4***"), |_, _span| Multicast::Warpx4)
61            )
62        }
63    }
64
65    impl PtxParser for Shape {
66        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
67            alt!(
68                map(string_p(".32x128b***"), |_, _span| Shape::_32x128b),
69                map(string_p(".64x128b**"), |_, _span| Shape::_64x128b),
70                map(string_p(".128x256b"), |_, _span| Shape::_128x256b),
71                map(string_p(".128x128b"), |_, _span| Shape::_128x128b),
72                map(string_p(".4x256b"), |_, _span| Shape::_4x256b)
73            )
74        }
75    }
76
77    impl PtxParser for Tcgen05CpCtaGroupShapeMulticastDstSrcFmt {
78        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
79            try_map(
80                seq_n!(
81                    string_p("tcgen05"),
82                    string_p(".cp"),
83                    CtaGroup::parse(),
84                    Shape::parse(),
85                    optional(Multicast::parse()),
86                    optional(DstSrcFmt::parse()),
87                    AddressOperand::parse(),
88                    comma_p(),
89                    GeneralOperand::parse(),
90                    semicolon_p()
91                ),
92                |(_, cp, cta_group, shape, multicast, dst_src_fmt, taddr, _, s_desc, _), span| {
93                    ok!(Tcgen05CpCtaGroupShapeMulticastDstSrcFmt {
94                        cp = cp,
95                        cta_group = cta_group,
96                        shape = shape,
97                        multicast = multicast,
98                        dst_src_fmt = dst_src_fmt,
99                        taddr = taddr,
100                        s_desc = s_desc,
101
102                    })
103                },
104            )
105        }
106    }
107}
ptx_parser/parser/instruction/tcgen05_cp.rs

ptx_parser/parser/instruction/
tcgen05_cp.rs
