ptx_parser/unparser/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::lexer::PtxToken;
12use crate::unparser::{PtxUnparser, common::*};
13
14pub mod section_0 {
15    use super::*;
16    use crate::r#type::instruction::tcgen05_cp::section_0::*;
17
18    impl PtxUnparser for Tcgen05CpCtaGroupShapeMulticastDstSrcFmt {
19        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
20            push_opcode(tokens, "tcgen05");
21            push_directive(tokens, "cp");
22            match &self.cta_group {
23                CtaGroup::CtaGroup1 => {
24                    push_directive(tokens, "cta_group::1");
25                }
26                CtaGroup::CtaGroup2 => {
27                    push_directive(tokens, "cta_group::2");
28                }
29            }
30            match &self.shape {
31                Shape::_32x128b => {
32                    push_directive(tokens, "32x128b***");
33                }
34                Shape::_64x128b => {
35                    push_directive(tokens, "64x128b**");
36                }
37                Shape::_128x256b => {
38                    push_directive(tokens, "128x256b");
39                }
40                Shape::_128x128b => {
41                    push_directive(tokens, "128x128b");
42                }
43                Shape::_4x256b => {
44                    push_directive(tokens, "4x256b");
45                }
46            }
47            if let Some(multicast_0) = self.multicast.as_ref() {
48                match multicast_0 {
49                    Multicast::Warpx20213 => {
50                        push_directive(tokens, "warpx2::02_13**");
51                    }
52                    Multicast::Warpx20123 => {
53                        push_directive(tokens, "warpx2::01_23**");
54                    }
55                    Multicast::Warpx4 => {
56                        push_directive(tokens, "warpx4***");
57                    }
58                }
59            }
60            if let Some(dst_src_fmt_1) = self.dst_src_fmt.as_ref() {
61                match dst_src_fmt_1 {
62                    DstSrcFmt::B8x16B6x16P32 => {
63                        push_directive(tokens, "b8x16.b6x16_p32");
64                    }
65                    DstSrcFmt::B8x16B4x16P64 => {
66                        push_directive(tokens, "b8x16.b4x16_p64");
67                    }
68                }
69            }
70            self.taddr.unparse_tokens(tokens);
71            tokens.push(PtxToken::Comma);
72            self.s_desc.unparse_tokens(tokens);
73            tokens.push(PtxToken::Semicolon);
74        }
75    }
76}
ptx_parser/unparser/instruction/tcgen05_cp.rs

ptx_parser/unparser/instruction/
tcgen05_cp.rs
