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}
ptx_parser/unparser/instruction/tcgen05_ld.rs

ptx_parser/unparser/instruction/
tcgen05_ld.rs
