ptx_parser/unparser/instruction/
ldmatrix.rs

1//! Original PTX specification:
2//!
3//! ldmatrix.sync.aligned.shape.num{.trans}{.ss}.type r, [p];
4//! ldmatrix.sync.aligned.m8n16.num{.ss}.dst_fmt.src_fmt        r, [p];
5//! ldmatrix.sync.aligned.m16n16.num.trans{.ss}.dst_fmt.src_fmt r, [p];
6//! .shape   = {.m8n8, .m16n16};
7//! .num     = {.x1, .x2, .x4};
8//! .ss      = {.shared, .shared::cta};
9//! .type    = {.b16, .b8};
10//! .dst_fmt = { .b8x16 };
11//! .src_fmt = { .b6x16_p32, .b4x16_p64 };
12
13#![allow(unused)]
14
15use crate::lexer::PtxToken;
16use crate::unparser::{PtxUnparser, common::*};
17
18pub mod section_0 {
19    use super::*;
20    use crate::r#type::instruction::ldmatrix::section_0::*;
21
22    impl PtxUnparser for LdmatrixSyncAlignedShapeNumTransSsType {
23        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
24            push_opcode(tokens, "ldmatrix");
25            push_directive(tokens, "sync");
26            push_directive(tokens, "aligned");
27            match &self.shape {
28                Shape::M16n16 => {
29                    push_directive(tokens, "m16n16");
30                }
31                Shape::M8n8 => {
32                    push_directive(tokens, "m8n8");
33                }
34            }
35            match &self.num {
36                Num::X1 => {
37                    push_directive(tokens, "x1");
38                }
39                Num::X2 => {
40                    push_directive(tokens, "x2");
41                }
42                Num::X4 => {
43                    push_directive(tokens, "x4");
44                }
45            }
46            if self.trans {
47                push_directive(tokens, "trans");
48            }
49            if let Some(ss_0) = self.ss.as_ref() {
50                match ss_0 {
51                    Ss::SharedCta => {
52                        push_directive(tokens, "shared::cta");
53                    }
54                    Ss::Shared => {
55                        push_directive(tokens, "shared");
56                    }
57                }
58            }
59            match &self.type_ {
60                Type::B16 => {
61                    push_directive(tokens, "b16");
62                }
63                Type::B8 => {
64                    push_directive(tokens, "b8");
65                }
66            }
67            self.r.unparse_tokens(tokens);
68            tokens.push(PtxToken::Comma);
69            self.p.unparse_tokens(tokens);
70            tokens.push(PtxToken::Semicolon);
71        }
72    }
73
74    impl PtxUnparser for LdmatrixSyncAlignedM8n16NumSsDstFmtSrcFmt {
75        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
76            push_opcode(tokens, "ldmatrix");
77            push_directive(tokens, "sync");
78            push_directive(tokens, "aligned");
79            push_directive(tokens, "m8n16");
80            match &self.num {
81                Num::X1 => {
82                    push_directive(tokens, "x1");
83                }
84                Num::X2 => {
85                    push_directive(tokens, "x2");
86                }
87                Num::X4 => {
88                    push_directive(tokens, "x4");
89                }
90            }
91            if let Some(ss_1) = self.ss.as_ref() {
92                match ss_1 {
93                    Ss::SharedCta => {
94                        push_directive(tokens, "shared::cta");
95                    }
96                    Ss::Shared => {
97                        push_directive(tokens, "shared");
98                    }
99                }
100            }
101            match &self.dst_fmt {
102                DstFmt::B8x16 => {
103                    push_directive(tokens, "b8x16");
104                }
105            }
106            match &self.src_fmt {
107                SrcFmt::B6x16P32 => {
108                    push_directive(tokens, "b6x16_p32");
109                }
110                SrcFmt::B4x16P64 => {
111                    push_directive(tokens, "b4x16_p64");
112                }
113            }
114            self.r.unparse_tokens(tokens);
115            tokens.push(PtxToken::Comma);
116            self.p.unparse_tokens(tokens);
117            tokens.push(PtxToken::Semicolon);
118        }
119    }
120
121    impl PtxUnparser for LdmatrixSyncAlignedM16n16NumTransSsDstFmtSrcFmt {
122        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
123            push_opcode(tokens, "ldmatrix");
124            push_directive(tokens, "sync");
125            push_directive(tokens, "aligned");
126            push_directive(tokens, "m16n16");
127            match &self.num {
128                Num::X1 => {
129                    push_directive(tokens, "x1");
130                }
131                Num::X2 => {
132                    push_directive(tokens, "x2");
133                }
134                Num::X4 => {
135                    push_directive(tokens, "x4");
136                }
137            }
138            push_directive(tokens, "trans");
139            if let Some(ss_2) = self.ss.as_ref() {
140                match ss_2 {
141                    Ss::SharedCta => {
142                        push_directive(tokens, "shared::cta");
143                    }
144                    Ss::Shared => {
145                        push_directive(tokens, "shared");
146                    }
147                }
148            }
149            match &self.dst_fmt {
150                DstFmt::B8x16 => {
151                    push_directive(tokens, "b8x16");
152                }
153            }
154            match &self.src_fmt {
155                SrcFmt::B6x16P32 => {
156                    push_directive(tokens, "b6x16_p32");
157                }
158                SrcFmt::B4x16P64 => {
159                    push_directive(tokens, "b4x16_p64");
160                }
161            }
162            self.r.unparse_tokens(tokens);
163            tokens.push(PtxToken::Comma);
164            self.p.unparse_tokens(tokens);
165            tokens.push(PtxToken::Semicolon);
166        }
167    }
168}
ptx_parser/unparser/instruction/ldmatrix.rs

ptx_parser/unparser/instruction/
ldmatrix.rs
