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
77}
78
ptx_parser/unparser/instruction/tcgen05_cp.rs

ptx_parser/unparser/instruction/
tcgen05_cp.rs
