ptx_parser/unparser/instruction/tcgen05_ld.rs
1//! Original PTX specification:
2//!
3//! // Base load instruction:
4//! tcgen05.ld.sync.aligned.shape1.num{.pack}.b32 r, [taddr];
5//! tcgen05.ld.sync.aligned.shape2.num{.pack}.b32 r, [taddr], immHalfSplitoff;
6//! .shape1 = { .16x64b, .16x128b, .16x256b, .32x32b };
7//! .shape2 = { .16x32bx2 };
8//! .num = { .x1, .x2, .x4, .x8, .x16, .x32, .x64, .x128 };
9//! .pack = { .pack::16b };
10//! // Floating point type load along with reduction :
11//! tcgen05.ld.red.sync.aligned.shape3.num.redOp{.abs}{.NaN}.f32 r, redval, [taddr];
12//! tcgen05.ld.red.sync.aligned.shape4.num.redOp{.abs}{.NaN}.f32 r, redval, [taddr], immHalfSplitoff;
13//! // Integer type load along with reduction :
14//! tcgen05.ld.red.sync.aligned.shape3.num.redOp.type r, redval, [taddr];
15//! tcgen05.ld.red.sync.aligned.shape4.num.redOp.type r, redval, [taddr], immHalfSplitoff;
16//! .shape3 = { .32x32b };
17//! .shape4 = { .16x32bx2 };
18//! .redOp = { .min, .max };
19//! .type = { .u32, .s32 };
20
21#![allow(unused)]
22
23use crate::lexer::PtxToken;
24use crate::unparser::{PtxUnparser, common::*};
25
26pub mod section_0 {
27 use super::*;
28 use crate::r#type::instruction::tcgen05_ld::section_0::*;
29
30 impl PtxUnparser for Tcgen05LdSyncAlignedShape1NumPackB32 {
31 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
32 push_opcode(tokens, "tcgen05");
33 push_directive(tokens, "ld");
34 push_directive(tokens, "sync");
35 push_directive(tokens, "aligned");
36 match &self.shape1 {
37 Shape1::_16x128b => {
38 push_directive(tokens, "16x128b");
39 }
40 Shape1::_16x256b => {
41 push_directive(tokens, "16x256b");
42 }
43 Shape1::_16x64b => {
44 push_directive(tokens, "16x64b");
45 }
46 Shape1::_32x32b => {
47 push_directive(tokens, "32x32b");
48 }
49 }
50 match &self.num {
51 Num::X128 => {
52 push_directive(tokens, "x128");
53 }
54 Num::X16 => {
55 push_directive(tokens, "x16");
56 }
57 Num::X32 => {
58 push_directive(tokens, "x32");
59 }
60 Num::X64 => {
61 push_directive(tokens, "x64");
62 }
63 Num::X1 => {
64 push_directive(tokens, "x1");
65 }
66 Num::X2 => {
67 push_directive(tokens, "x2");
68 }
69 Num::X4 => {
70 push_directive(tokens, "x4");
71 }
72 Num::X8 => {
73 push_directive(tokens, "x8");
74 }
75 }
76 if let Some(pack_0) = self.pack.as_ref() {
77 match pack_0 {
78 Pack::Pack16b => {
79 push_directive(tokens, "pack::16b");
80 }
81 }
82 }
83 push_directive(tokens, "b32");
84 self.r.unparse_tokens(tokens);
85 tokens.push(PtxToken::Comma);
86 self.taddr.unparse_tokens(tokens);
87 tokens.push(PtxToken::Semicolon);
88 }
89 }
90
91 impl PtxUnparser for Tcgen05LdSyncAlignedShape2NumPackB32 {
92 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
93 push_opcode(tokens, "tcgen05");
94 push_directive(tokens, "ld");
95 push_directive(tokens, "sync");
96 push_directive(tokens, "aligned");
97 match &self.shape2 {
98 Shape2::_16x32bx2 => {
99 push_directive(tokens, "16x32bx2");
100 }
101 }
102 match &self.num {
103 Num::X128 => {
104 push_directive(tokens, "x128");
105 }
106 Num::X16 => {
107 push_directive(tokens, "x16");
108 }
109 Num::X32 => {
110 push_directive(tokens, "x32");
111 }
112 Num::X64 => {
113 push_directive(tokens, "x64");
114 }
115 Num::X1 => {
116 push_directive(tokens, "x1");
117 }
118 Num::X2 => {
119 push_directive(tokens, "x2");
120 }
121 Num::X4 => {
122 push_directive(tokens, "x4");
123 }
124 Num::X8 => {
125 push_directive(tokens, "x8");
126 }
127 }
128 if let Some(pack_1) = self.pack.as_ref() {
129 match pack_1 {
130 Pack::Pack16b => {
131 push_directive(tokens, "pack::16b");
132 }
133 }
134 }
135 push_directive(tokens, "b32");
136 self.r.unparse_tokens(tokens);
137 tokens.push(PtxToken::Comma);
138 self.taddr.unparse_tokens(tokens);
139 tokens.push(PtxToken::Comma);
140 self.immhalfsplitoff.unparse_tokens(tokens);
141 tokens.push(PtxToken::Semicolon);
142 }
143 }
144
145 impl PtxUnparser for Tcgen05LdRedSyncAlignedShape3NumRedopAbsNanF32 {
146 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
147 push_opcode(tokens, "tcgen05");
148 push_directive(tokens, "ld");
149 push_directive(tokens, "red");
150 push_directive(tokens, "sync");
151 push_directive(tokens, "aligned");
152 match &self.shape3 {
153 Shape3::_32x32b => {
154 push_directive(tokens, "32x32b");
155 }
156 }
157 match &self.num {
158 Num::X128 => {
159 push_directive(tokens, "x128");
160 }
161 Num::X16 => {
162 push_directive(tokens, "x16");
163 }
164 Num::X32 => {
165 push_directive(tokens, "x32");
166 }
167 Num::X64 => {
168 push_directive(tokens, "x64");
169 }
170 Num::X1 => {
171 push_directive(tokens, "x1");
172 }
173 Num::X2 => {
174 push_directive(tokens, "x2");
175 }
176 Num::X4 => {
177 push_directive(tokens, "x4");
178 }
179 Num::X8 => {
180 push_directive(tokens, "x8");
181 }
182 }
183 match &self.redop {
184 Redop::Min => {
185 push_directive(tokens, "min");
186 }
187 Redop::Max => {
188 push_directive(tokens, "max");
189 }
190 }
191 if self.abs {
192 push_directive(tokens, "abs");
193 }
194 if self.nan {
195 push_directive(tokens, "NaN");
196 }
197 push_directive(tokens, "f32");
198 self.r.unparse_tokens(tokens);
199 tokens.push(PtxToken::Comma);
200 self.redval.unparse_tokens(tokens);
201 tokens.push(PtxToken::Comma);
202 self.taddr.unparse_tokens(tokens);
203 tokens.push(PtxToken::Semicolon);
204 }
205 }
206
207 impl PtxUnparser for Tcgen05LdRedSyncAlignedShape4NumRedopAbsNanF32 {
208 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
209 push_opcode(tokens, "tcgen05");
210 push_directive(tokens, "ld");
211 push_directive(tokens, "red");
212 push_directive(tokens, "sync");
213 push_directive(tokens, "aligned");
214 match &self.shape4 {
215 Shape4::_16x32bx2 => {
216 push_directive(tokens, "16x32bx2");
217 }
218 }
219 match &self.num {
220 Num::X128 => {
221 push_directive(tokens, "x128");
222 }
223 Num::X16 => {
224 push_directive(tokens, "x16");
225 }
226 Num::X32 => {
227 push_directive(tokens, "x32");
228 }
229 Num::X64 => {
230 push_directive(tokens, "x64");
231 }
232 Num::X1 => {
233 push_directive(tokens, "x1");
234 }
235 Num::X2 => {
236 push_directive(tokens, "x2");
237 }
238 Num::X4 => {
239 push_directive(tokens, "x4");
240 }
241 Num::X8 => {
242 push_directive(tokens, "x8");
243 }
244 }
245 match &self.redop {
246 Redop::Min => {
247 push_directive(tokens, "min");
248 }
249 Redop::Max => {
250 push_directive(tokens, "max");
251 }
252 }
253 if self.abs {
254 push_directive(tokens, "abs");
255 }
256 if self.nan {
257 push_directive(tokens, "NaN");
258 }
259 push_directive(tokens, "f32");
260 self.r.unparse_tokens(tokens);
261 tokens.push(PtxToken::Comma);
262 self.redval.unparse_tokens(tokens);
263 tokens.push(PtxToken::Comma);
264 self.taddr.unparse_tokens(tokens);
265 tokens.push(PtxToken::Comma);
266 self.immhalfsplitoff.unparse_tokens(tokens);
267 tokens.push(PtxToken::Semicolon);
268 }
269 }
270
271 impl PtxUnparser for Tcgen05LdRedSyncAlignedShape3NumRedopType {
272 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
273 push_opcode(tokens, "tcgen05");
274 push_directive(tokens, "ld");
275 push_directive(tokens, "red");
276 push_directive(tokens, "sync");
277 push_directive(tokens, "aligned");
278 match &self.shape3 {
279 Shape3::_32x32b => {
280 push_directive(tokens, "32x32b");
281 }
282 }
283 match &self.num {
284 Num::X128 => {
285 push_directive(tokens, "x128");
286 }
287 Num::X16 => {
288 push_directive(tokens, "x16");
289 }
290 Num::X32 => {
291 push_directive(tokens, "x32");
292 }
293 Num::X64 => {
294 push_directive(tokens, "x64");
295 }
296 Num::X1 => {
297 push_directive(tokens, "x1");
298 }
299 Num::X2 => {
300 push_directive(tokens, "x2");
301 }
302 Num::X4 => {
303 push_directive(tokens, "x4");
304 }
305 Num::X8 => {
306 push_directive(tokens, "x8");
307 }
308 }
309 match &self.redop {
310 Redop::Min => {
311 push_directive(tokens, "min");
312 }
313 Redop::Max => {
314 push_directive(tokens, "max");
315 }
316 }
317 match &self.type_ {
318 Type::U32 => {
319 push_directive(tokens, "u32");
320 }
321 Type::S32 => {
322 push_directive(tokens, "s32");
323 }
324 }
325 self.r.unparse_tokens(tokens);
326 tokens.push(PtxToken::Comma);
327 self.redval.unparse_tokens(tokens);
328 tokens.push(PtxToken::Comma);
329 self.taddr.unparse_tokens(tokens);
330 tokens.push(PtxToken::Semicolon);
331 }
332 }
333
334 impl PtxUnparser for Tcgen05LdRedSyncAlignedShape4NumRedopType {
335 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
336 push_opcode(tokens, "tcgen05");
337 push_directive(tokens, "ld");
338 push_directive(tokens, "red");
339 push_directive(tokens, "sync");
340 push_directive(tokens, "aligned");
341 match &self.shape4 {
342 Shape4::_16x32bx2 => {
343 push_directive(tokens, "16x32bx2");
344 }
345 }
346 match &self.num {
347 Num::X128 => {
348 push_directive(tokens, "x128");
349 }
350 Num::X16 => {
351 push_directive(tokens, "x16");
352 }
353 Num::X32 => {
354 push_directive(tokens, "x32");
355 }
356 Num::X64 => {
357 push_directive(tokens, "x64");
358 }
359 Num::X1 => {
360 push_directive(tokens, "x1");
361 }
362 Num::X2 => {
363 push_directive(tokens, "x2");
364 }
365 Num::X4 => {
366 push_directive(tokens, "x4");
367 }
368 Num::X8 => {
369 push_directive(tokens, "x8");
370 }
371 }
372 match &self.redop {
373 Redop::Min => {
374 push_directive(tokens, "min");
375 }
376 Redop::Max => {
377 push_directive(tokens, "max");
378 }
379 }
380 match &self.type_ {
381 Type::U32 => {
382 push_directive(tokens, "u32");
383 }
384 Type::S32 => {
385 push_directive(tokens, "s32");
386 }
387 }
388 self.r.unparse_tokens(tokens);
389 tokens.push(PtxToken::Comma);
390 self.redval.unparse_tokens(tokens);
391 tokens.push(PtxToken::Comma);
392 self.taddr.unparse_tokens(tokens);
393 tokens.push(PtxToken::Comma);
394 self.immhalfsplitoff.unparse_tokens(tokens);
395 tokens.push(PtxToken::Semicolon);
396 }
397 }
398
399}
400