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            self.unparse_tokens_mode(tokens, false);
33        }
34        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
35            push_opcode(tokens, "tcgen05");
36            push_directive(tokens, "ld");
37            push_directive(tokens, "sync");
38            push_directive(tokens, "aligned");
39            match &self.shape1 {
40                Shape1::_16x128b => {
41                    push_directive(tokens, "16x128b");
42                }
43                Shape1::_16x256b => {
44                    push_directive(tokens, "16x256b");
45                }
46                Shape1::_16x64b => {
47                    push_directive(tokens, "16x64b");
48                }
49                Shape1::_32x32b => {
50                    push_directive(tokens, "32x32b");
51                }
52            }
53            match &self.num {
54                Num::X128 => {
55                    push_directive(tokens, "x128");
56                }
57                Num::X16 => {
58                    push_directive(tokens, "x16");
59                }
60                Num::X32 => {
61                    push_directive(tokens, "x32");
62                }
63                Num::X64 => {
64                    push_directive(tokens, "x64");
65                }
66                Num::X1 => {
67                    push_directive(tokens, "x1");
68                }
69                Num::X2 => {
70                    push_directive(tokens, "x2");
71                }
72                Num::X4 => {
73                    push_directive(tokens, "x4");
74                }
75                Num::X8 => {
76                    push_directive(tokens, "x8");
77                }
78            }
79            if let Some(pack_0) = self.pack.as_ref() {
80                match pack_0 {
81                    Pack::Pack16b => {
82                        push_directive(tokens, "pack::16b");
83                    }
84                }
85            }
86            push_directive(tokens, "b32");
87            if spaced {
88                tokens.push(PtxToken::Space);
89            }
90            self.r.unparse_tokens_mode(tokens, spaced);
91            tokens.push(PtxToken::Comma);
92            if spaced {
93                tokens.push(PtxToken::Space);
94            }
95            self.taddr.unparse_tokens_mode(tokens, spaced);
96            tokens.push(PtxToken::Semicolon);
97            if spaced {
98                tokens.push(PtxToken::Newline);
99            }
100        }
101    }
102
103    impl PtxUnparser for Tcgen05LdSyncAlignedShape2NumPackB32 {
104        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
105            self.unparse_tokens_mode(tokens, false);
106        }
107        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
108            push_opcode(tokens, "tcgen05");
109            push_directive(tokens, "ld");
110            push_directive(tokens, "sync");
111            push_directive(tokens, "aligned");
112            match &self.shape2 {
113                Shape2::_16x32bx2 => {
114                    push_directive(tokens, "16x32bx2");
115                }
116            }
117            match &self.num {
118                Num::X128 => {
119                    push_directive(tokens, "x128");
120                }
121                Num::X16 => {
122                    push_directive(tokens, "x16");
123                }
124                Num::X32 => {
125                    push_directive(tokens, "x32");
126                }
127                Num::X64 => {
128                    push_directive(tokens, "x64");
129                }
130                Num::X1 => {
131                    push_directive(tokens, "x1");
132                }
133                Num::X2 => {
134                    push_directive(tokens, "x2");
135                }
136                Num::X4 => {
137                    push_directive(tokens, "x4");
138                }
139                Num::X8 => {
140                    push_directive(tokens, "x8");
141                }
142            }
143            if let Some(pack_1) = self.pack.as_ref() {
144                match pack_1 {
145                    Pack::Pack16b => {
146                        push_directive(tokens, "pack::16b");
147                    }
148                }
149            }
150            push_directive(tokens, "b32");
151            if spaced {
152                tokens.push(PtxToken::Space);
153            }
154            self.r.unparse_tokens_mode(tokens, spaced);
155            tokens.push(PtxToken::Comma);
156            if spaced {
157                tokens.push(PtxToken::Space);
158            }
159            self.taddr.unparse_tokens_mode(tokens, spaced);
160            tokens.push(PtxToken::Comma);
161            if spaced {
162                tokens.push(PtxToken::Space);
163            }
164            self.immhalfsplitoff.unparse_tokens_mode(tokens, spaced);
165            tokens.push(PtxToken::Semicolon);
166            if spaced {
167                tokens.push(PtxToken::Newline);
168            }
169        }
170    }
171
172    impl PtxUnparser for Tcgen05LdRedSyncAlignedShape3NumRedopAbsNanF32 {
173        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
174            self.unparse_tokens_mode(tokens, false);
175        }
176        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
177            push_opcode(tokens, "tcgen05");
178            push_directive(tokens, "ld");
179            push_directive(tokens, "red");
180            push_directive(tokens, "sync");
181            push_directive(tokens, "aligned");
182            match &self.shape3 {
183                Shape3::_32x32b => {
184                    push_directive(tokens, "32x32b");
185                }
186            }
187            match &self.num {
188                Num::X128 => {
189                    push_directive(tokens, "x128");
190                }
191                Num::X16 => {
192                    push_directive(tokens, "x16");
193                }
194                Num::X32 => {
195                    push_directive(tokens, "x32");
196                }
197                Num::X64 => {
198                    push_directive(tokens, "x64");
199                }
200                Num::X1 => {
201                    push_directive(tokens, "x1");
202                }
203                Num::X2 => {
204                    push_directive(tokens, "x2");
205                }
206                Num::X4 => {
207                    push_directive(tokens, "x4");
208                }
209                Num::X8 => {
210                    push_directive(tokens, "x8");
211                }
212            }
213            match &self.redop {
214                Redop::Min => {
215                    push_directive(tokens, "min");
216                }
217                Redop::Max => {
218                    push_directive(tokens, "max");
219                }
220            }
221            if self.abs {
222                push_directive(tokens, "abs");
223            }
224            if self.nan {
225                push_directive(tokens, "NaN");
226            }
227            push_directive(tokens, "f32");
228            if spaced {
229                tokens.push(PtxToken::Space);
230            }
231            self.r.unparse_tokens_mode(tokens, spaced);
232            tokens.push(PtxToken::Comma);
233            if spaced {
234                tokens.push(PtxToken::Space);
235            }
236            self.redval.unparse_tokens_mode(tokens, spaced);
237            tokens.push(PtxToken::Comma);
238            if spaced {
239                tokens.push(PtxToken::Space);
240            }
241            self.taddr.unparse_tokens_mode(tokens, spaced);
242            tokens.push(PtxToken::Semicolon);
243            if spaced {
244                tokens.push(PtxToken::Newline);
245            }
246        }
247    }
248
249    impl PtxUnparser for Tcgen05LdRedSyncAlignedShape4NumRedopAbsNanF32 {
250        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
251            self.unparse_tokens_mode(tokens, false);
252        }
253        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
254            push_opcode(tokens, "tcgen05");
255            push_directive(tokens, "ld");
256            push_directive(tokens, "red");
257            push_directive(tokens, "sync");
258            push_directive(tokens, "aligned");
259            match &self.shape4 {
260                Shape4::_16x32bx2 => {
261                    push_directive(tokens, "16x32bx2");
262                }
263            }
264            match &self.num {
265                Num::X128 => {
266                    push_directive(tokens, "x128");
267                }
268                Num::X16 => {
269                    push_directive(tokens, "x16");
270                }
271                Num::X32 => {
272                    push_directive(tokens, "x32");
273                }
274                Num::X64 => {
275                    push_directive(tokens, "x64");
276                }
277                Num::X1 => {
278                    push_directive(tokens, "x1");
279                }
280                Num::X2 => {
281                    push_directive(tokens, "x2");
282                }
283                Num::X4 => {
284                    push_directive(tokens, "x4");
285                }
286                Num::X8 => {
287                    push_directive(tokens, "x8");
288                }
289            }
290            match &self.redop {
291                Redop::Min => {
292                    push_directive(tokens, "min");
293                }
294                Redop::Max => {
295                    push_directive(tokens, "max");
296                }
297            }
298            if self.abs {
299                push_directive(tokens, "abs");
300            }
301            if self.nan {
302                push_directive(tokens, "NaN");
303            }
304            push_directive(tokens, "f32");
305            if spaced {
306                tokens.push(PtxToken::Space);
307            }
308            self.r.unparse_tokens_mode(tokens, spaced);
309            tokens.push(PtxToken::Comma);
310            if spaced {
311                tokens.push(PtxToken::Space);
312            }
313            self.redval.unparse_tokens_mode(tokens, spaced);
314            tokens.push(PtxToken::Comma);
315            if spaced {
316                tokens.push(PtxToken::Space);
317            }
318            self.taddr.unparse_tokens_mode(tokens, spaced);
319            tokens.push(PtxToken::Comma);
320            if spaced {
321                tokens.push(PtxToken::Space);
322            }
323            self.immhalfsplitoff.unparse_tokens_mode(tokens, spaced);
324            tokens.push(PtxToken::Semicolon);
325            if spaced {
326                tokens.push(PtxToken::Newline);
327            }
328        }
329    }
330
331    impl PtxUnparser for Tcgen05LdRedSyncAlignedShape3NumRedopType {
332        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
333            self.unparse_tokens_mode(tokens, false);
334        }
335        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
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.shape3 {
342                Shape3::_32x32b => {
343                    push_directive(tokens, "32x32b");
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            if spaced {
389                tokens.push(PtxToken::Space);
390            }
391            self.r.unparse_tokens_mode(tokens, spaced);
392            tokens.push(PtxToken::Comma);
393            if spaced {
394                tokens.push(PtxToken::Space);
395            }
396            self.redval.unparse_tokens_mode(tokens, spaced);
397            tokens.push(PtxToken::Comma);
398            if spaced {
399                tokens.push(PtxToken::Space);
400            }
401            self.taddr.unparse_tokens_mode(tokens, spaced);
402            tokens.push(PtxToken::Semicolon);
403            if spaced {
404                tokens.push(PtxToken::Newline);
405            }
406        }
407    }
408
409    impl PtxUnparser for Tcgen05LdRedSyncAlignedShape4NumRedopType {
410        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
411            self.unparse_tokens_mode(tokens, false);
412        }
413        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
414            push_opcode(tokens, "tcgen05");
415            push_directive(tokens, "ld");
416            push_directive(tokens, "red");
417            push_directive(tokens, "sync");
418            push_directive(tokens, "aligned");
419            match &self.shape4 {
420                Shape4::_16x32bx2 => {
421                    push_directive(tokens, "16x32bx2");
422                }
423            }
424            match &self.num {
425                Num::X128 => {
426                    push_directive(tokens, "x128");
427                }
428                Num::X16 => {
429                    push_directive(tokens, "x16");
430                }
431                Num::X32 => {
432                    push_directive(tokens, "x32");
433                }
434                Num::X64 => {
435                    push_directive(tokens, "x64");
436                }
437                Num::X1 => {
438                    push_directive(tokens, "x1");
439                }
440                Num::X2 => {
441                    push_directive(tokens, "x2");
442                }
443                Num::X4 => {
444                    push_directive(tokens, "x4");
445                }
446                Num::X8 => {
447                    push_directive(tokens, "x8");
448                }
449            }
450            match &self.redop {
451                Redop::Min => {
452                    push_directive(tokens, "min");
453                }
454                Redop::Max => {
455                    push_directive(tokens, "max");
456                }
457            }
458            match &self.type_ {
459                Type::U32 => {
460                    push_directive(tokens, "u32");
461                }
462                Type::S32 => {
463                    push_directive(tokens, "s32");
464                }
465            }
466            if spaced {
467                tokens.push(PtxToken::Space);
468            }
469            self.r.unparse_tokens_mode(tokens, spaced);
470            tokens.push(PtxToken::Comma);
471            if spaced {
472                tokens.push(PtxToken::Space);
473            }
474            self.redval.unparse_tokens_mode(tokens, spaced);
475            tokens.push(PtxToken::Comma);
476            if spaced {
477                tokens.push(PtxToken::Space);
478            }
479            self.taddr.unparse_tokens_mode(tokens, spaced);
480            tokens.push(PtxToken::Comma);
481            if spaced {
482                tokens.push(PtxToken::Space);
483            }
484            self.immhalfsplitoff.unparse_tokens_mode(tokens, spaced);
485            tokens.push(PtxToken::Semicolon);
486            if spaced {
487                tokens.push(PtxToken::Newline);
488            }
489        }
490    }
491}
ptx_parser/unparser/instruction/tcgen05_ld.rs

ptx_parser/unparser/instruction/
tcgen05_ld.rs
