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            self.unparse_tokens_mode(tokens, false);
21        }
22        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
23            push_opcode(tokens, "tcgen05");
24            push_directive(tokens, "cp");
25            match &self.cta_group {
26                CtaGroup::CtaGroup1 => {
27                    push_directive(tokens, "cta_group::1");
28                }
29                CtaGroup::CtaGroup2 => {
30                    push_directive(tokens, "cta_group::2");
31                }
32            }
33            match &self.shape {
34                Shape::_32x128b => {
35                    push_directive(tokens, "32x128b***");
36                }
37                Shape::_64x128b => {
38                    push_directive(tokens, "64x128b**");
39                }
40                Shape::_128x256b => {
41                    push_directive(tokens, "128x256b");
42                }
43                Shape::_128x128b => {
44                    push_directive(tokens, "128x128b");
45                }
46                Shape::_4x256b => {
47                    push_directive(tokens, "4x256b");
48                }
49            }
50            if let Some(multicast_0) = self.multicast.as_ref() {
51                match multicast_0 {
52                    Multicast::Warpx20213 => {
53                        push_directive(tokens, "warpx2::02_13**");
54                    }
55                    Multicast::Warpx20123 => {
56                        push_directive(tokens, "warpx2::01_23**");
57                    }
58                    Multicast::Warpx4 => {
59                        push_directive(tokens, "warpx4***");
60                    }
61                }
62            }
63            if let Some(dst_src_fmt_1) = self.dst_src_fmt.as_ref() {
64                match dst_src_fmt_1 {
65                    DstSrcFmt::B8x16B6x16P32 => {
66                        push_directive(tokens, "b8x16.b6x16_p32");
67                    }
68                    DstSrcFmt::B8x16B4x16P64 => {
69                        push_directive(tokens, "b8x16.b4x16_p64");
70                    }
71                }
72            }
73            if spaced {
74                tokens.push(PtxToken::Space);
75            }
76            self.taddr.unparse_tokens_mode(tokens, spaced);
77            tokens.push(PtxToken::Comma);
78            if spaced {
79                tokens.push(PtxToken::Space);
80            }
81            self.s_desc.unparse_tokens_mode(tokens, spaced);
82            tokens.push(PtxToken::Semicolon);
83            if spaced {
84                tokens.push(PtxToken::Newline);
85            }
86        }
87    }
88}
ptx_parser/unparser/instruction/tcgen05_cp.rs

ptx_parser/unparser/instruction/
tcgen05_cp.rs
