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
169}
170