ptx_parser/unparser/instruction/
cp_async_bulk_tensor.rs

1//! Original PTX specification:
2//!
3//! // global -> shared::cta
4//! cp.async.bulk.tensor.dim.dst.src{.load_mode}.completion_mechanism{.cta_group}{.level::cache_hint} [dstMem], [tensorMap, tensorCoords], [mbar]{, im2colInfo} {, cache-policy};
5//! .dst =                  { .shared::cta };
6//! .src =                  { .global };
7//! .dim =                  { .1d, .2d, .3d, .4d, .5d };
8//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
9//! .cta_group =            { .cta_group::1, .cta_group::2 };
10//! .load_mode =            { .tile, .tile::gather4, .im2col, .im2col::w, .im2col::w::128 };
11//! .level::cache_hint =    { .L2::cache_hint };
12//! ----------------------------------------------------------------
13//! // global -> shared::cluster
14//! cp.async.bulk.tensor.dim.dst.src{.load_mode}.completion_mechanism{.multicast}{.cta_group}{.level::cache_hint} [dstMem], [tensorMap, tensorCoords], [mbar]{, im2colInfo} {, ctaMask} {, cache-policy};
15//! .dst =                  { .shared::cluster };
16//! .src =                  { .global };
17//! .dim =                  { .1d, .2d, .3d, .4d, .5d };
18//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
19//! .cta_group =            { .cta_group::1, .cta_group::2 };
20//! .load_mode =            { .tile, .tile::gather4, .im2col, .im2col::w, .im2col::w::128 };
21//! .level::cache_hint =    { .L2::cache_hint };
22//! .multicast =            { .multicast::cluster  };
23//! ----------------------------------------------------------------
24//! // shared::cta -> global;
25//! cp.async.bulk.tensor.dim.dst.src{.load_mode}.completion_mechanism{.level::cache_hint} [tensorMap, tensorCoords], [srcMem] {, cache-policy};
26//! .dst =                  { .global };
27//! .src =                  { .shared::cta };
28//! .dim =                  { .1d, .2d, .3d, .4d, .5d };
29//! .completion_mechanism = { .bulk_group };
30//! .load_mode =            { .tile, .tile::scatter4, .im2col_no_offs };
31//! .level::cache_hint =    { .L2::cache_hint };
32
33#![allow(unused)]
34
35use crate::lexer::PtxToken;
36use crate::unparser::{PtxUnparser, common::*};
37
38pub mod section_0 {
39    use super::*;
40    use crate::r#type::instruction::cp_async_bulk_tensor::section_0::*;
41
42    impl PtxUnparser for CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismCtaGroupLevelCacheHint {
43        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
44            self.unparse_tokens_mode(tokens, false);
45        }
46        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
47            push_opcode(tokens, "cp");
48            push_directive(tokens, "async");
49            push_directive(tokens, "bulk");
50            push_directive(tokens, "tensor");
51            match &self.dim {
52                Dim::_1d => {
53                    push_directive(tokens, "1d");
54                }
55                Dim::_2d => {
56                    push_directive(tokens, "2d");
57                }
58                Dim::_3d => {
59                    push_directive(tokens, "3d");
60                }
61                Dim::_4d => {
62                    push_directive(tokens, "4d");
63                }
64                Dim::_5d => {
65                    push_directive(tokens, "5d");
66                }
67            }
68            match &self.dst {
69                Dst::SharedCta => {
70                    push_directive(tokens, "shared::cta");
71                }
72            }
73            match &self.src {
74                Src::Global => {
75                    push_directive(tokens, "global");
76                }
77            }
78            if let Some(load_mode_0) = self.load_mode.as_ref() {
79                match load_mode_0 {
80                    LoadMode::Im2colW128 => {
81                        push_directive(tokens, "im2col::w::128");
82                    }
83                    LoadMode::TileGather4 => {
84                        push_directive(tokens, "tile::gather4");
85                    }
86                    LoadMode::Im2colW => {
87                        push_directive(tokens, "im2col::w");
88                    }
89                    LoadMode::Im2col => {
90                        push_directive(tokens, "im2col");
91                    }
92                    LoadMode::Tile => {
93                        push_directive(tokens, "tile");
94                    }
95                }
96            }
97            match &self.completion_mechanism {
98                CompletionMechanism::MbarrierCompleteTxBytes => {
99                    push_directive(tokens, "mbarrier::complete_tx::bytes");
100                }
101            }
102            if let Some(cta_group_1) = self.cta_group.as_ref() {
103                match cta_group_1 {
104                    CtaGroup::CtaGroup1 => {
105                        push_directive(tokens, "cta_group::1");
106                    }
107                    CtaGroup::CtaGroup2 => {
108                        push_directive(tokens, "cta_group::2");
109                    }
110                }
111            }
112            if let Some(level_cache_hint_2) = self.level_cache_hint.as_ref() {
113                match level_cache_hint_2 {
114                    LevelCacheHint::L2CacheHint => {
115                        push_directive(tokens, "L2::cache_hint");
116                    }
117                }
118            }
119            if spaced {
120                tokens.push(PtxToken::Space);
121            }
122            self.dstmem.unparse_tokens_mode(tokens, spaced);
123            tokens.push(PtxToken::Comma);
124            if spaced {
125                tokens.push(PtxToken::Space);
126            }
127            self.tensormap.unparse_tokens_mode(tokens, spaced);
128            tokens.push(PtxToken::Comma);
129            if spaced {
130                tokens.push(PtxToken::Space);
131            }
132            self.mbar.unparse_tokens_mode(tokens, spaced);
133            if self.im2colinfo.is_some() {
134                tokens.push(PtxToken::Comma);
135            }
136            if let Some(opt_3) = self.im2colinfo.as_ref() {
137                if spaced {
138                    tokens.push(PtxToken::Space);
139                }
140                opt_3.unparse_tokens_mode(tokens, spaced);
141            }
142            if self.cache_policy.is_some() {
143                tokens.push(PtxToken::Comma);
144            }
145            if let Some(opt_4) = self.cache_policy.as_ref() {
146                if spaced {
147                    tokens.push(PtxToken::Space);
148                }
149                opt_4.unparse_tokens_mode(tokens, spaced);
150            }
151            tokens.push(PtxToken::Semicolon);
152            if spaced {
153                tokens.push(PtxToken::Newline);
154            }
155        }
156    }
157}
158
159pub mod section_1 {
160    use super::*;
161    use crate::r#type::instruction::cp_async_bulk_tensor::section_1::*;
162
163    impl PtxUnparser
164        for CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismMulticastCtaGroupLevelCacheHint
165    {
166        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
167            self.unparse_tokens_mode(tokens, false);
168        }
169        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
170            push_opcode(tokens, "cp");
171            push_directive(tokens, "async");
172            push_directive(tokens, "bulk");
173            push_directive(tokens, "tensor");
174            match &self.dim {
175                Dim::_1d => {
176                    push_directive(tokens, "1d");
177                }
178                Dim::_2d => {
179                    push_directive(tokens, "2d");
180                }
181                Dim::_3d => {
182                    push_directive(tokens, "3d");
183                }
184                Dim::_4d => {
185                    push_directive(tokens, "4d");
186                }
187                Dim::_5d => {
188                    push_directive(tokens, "5d");
189                }
190            }
191            match &self.dst {
192                Dst::SharedCluster => {
193                    push_directive(tokens, "shared::cluster");
194                }
195            }
196            match &self.src {
197                Src::Global => {
198                    push_directive(tokens, "global");
199                }
200            }
201            if let Some(load_mode_5) = self.load_mode.as_ref() {
202                match load_mode_5 {
203                    LoadMode::Im2colW128 => {
204                        push_directive(tokens, "im2col::w::128");
205                    }
206                    LoadMode::TileGather4 => {
207                        push_directive(tokens, "tile::gather4");
208                    }
209                    LoadMode::Im2colW => {
210                        push_directive(tokens, "im2col::w");
211                    }
212                    LoadMode::Im2col => {
213                        push_directive(tokens, "im2col");
214                    }
215                    LoadMode::Tile => {
216                        push_directive(tokens, "tile");
217                    }
218                }
219            }
220            match &self.completion_mechanism {
221                CompletionMechanism::MbarrierCompleteTxBytes => {
222                    push_directive(tokens, "mbarrier::complete_tx::bytes");
223                }
224            }
225            if let Some(multicast_6) = self.multicast.as_ref() {
226                match multicast_6 {
227                    Multicast::MulticastCluster => {
228                        push_directive(tokens, "multicast::cluster");
229                    }
230                }
231            }
232            if let Some(cta_group_7) = self.cta_group.as_ref() {
233                match cta_group_7 {
234                    CtaGroup::CtaGroup1 => {
235                        push_directive(tokens, "cta_group::1");
236                    }
237                    CtaGroup::CtaGroup2 => {
238                        push_directive(tokens, "cta_group::2");
239                    }
240                }
241            }
242            if let Some(level_cache_hint_8) = self.level_cache_hint.as_ref() {
243                match level_cache_hint_8 {
244                    LevelCacheHint::L2CacheHint => {
245                        push_directive(tokens, "L2::cache_hint");
246                    }
247                }
248            }
249            if spaced {
250                tokens.push(PtxToken::Space);
251            }
252            self.dstmem.unparse_tokens_mode(tokens, spaced);
253            tokens.push(PtxToken::Comma);
254            if spaced {
255                tokens.push(PtxToken::Space);
256            }
257            self.tensormap.unparse_tokens_mode(tokens, spaced);
258            tokens.push(PtxToken::Comma);
259            if spaced {
260                tokens.push(PtxToken::Space);
261            }
262            self.mbar.unparse_tokens_mode(tokens, spaced);
263            if self.im2colinfo.is_some() {
264                tokens.push(PtxToken::Comma);
265            }
266            if let Some(opt_9) = self.im2colinfo.as_ref() {
267                if spaced {
268                    tokens.push(PtxToken::Space);
269                }
270                opt_9.unparse_tokens_mode(tokens, spaced);
271            }
272            if self.ctamask.is_some() {
273                tokens.push(PtxToken::Comma);
274            }
275            if let Some(opt_10) = self.ctamask.as_ref() {
276                if spaced {
277                    tokens.push(PtxToken::Space);
278                }
279                opt_10.unparse_tokens_mode(tokens, spaced);
280            }
281            if self.cache_policy.is_some() {
282                tokens.push(PtxToken::Comma);
283            }
284            if let Some(opt_11) = self.cache_policy.as_ref() {
285                if spaced {
286                    tokens.push(PtxToken::Space);
287                }
288                opt_11.unparse_tokens_mode(tokens, spaced);
289            }
290            tokens.push(PtxToken::Semicolon);
291            if spaced {
292                tokens.push(PtxToken::Newline);
293            }
294        }
295    }
296}
297
298pub mod section_2 {
299    use super::*;
300    use crate::r#type::instruction::cp_async_bulk_tensor::section_2::*;
301
302    impl PtxUnparser for CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismLevelCacheHint {
303        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
304            self.unparse_tokens_mode(tokens, false);
305        }
306        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
307            push_opcode(tokens, "cp");
308            push_directive(tokens, "async");
309            push_directive(tokens, "bulk");
310            push_directive(tokens, "tensor");
311            match &self.dim {
312                Dim::_1d => {
313                    push_directive(tokens, "1d");
314                }
315                Dim::_2d => {
316                    push_directive(tokens, "2d");
317                }
318                Dim::_3d => {
319                    push_directive(tokens, "3d");
320                }
321                Dim::_4d => {
322                    push_directive(tokens, "4d");
323                }
324                Dim::_5d => {
325                    push_directive(tokens, "5d");
326                }
327            }
328            match &self.dst {
329                Dst::Global => {
330                    push_directive(tokens, "global");
331                }
332            }
333            match &self.src {
334                Src::SharedCta => {
335                    push_directive(tokens, "shared::cta");
336                }
337            }
338            if let Some(load_mode_12) = self.load_mode.as_ref() {
339                match load_mode_12 {
340                    LoadMode::TileScatter4 => {
341                        push_directive(tokens, "tile::scatter4");
342                    }
343                    LoadMode::Im2colNoOffs => {
344                        push_directive(tokens, "im2col_no_offs");
345                    }
346                    LoadMode::Tile => {
347                        push_directive(tokens, "tile");
348                    }
349                }
350            }
351            match &self.completion_mechanism {
352                CompletionMechanism::BulkGroup => {
353                    push_directive(tokens, "bulk_group");
354                }
355            }
356            if let Some(level_cache_hint_13) = self.level_cache_hint.as_ref() {
357                match level_cache_hint_13 {
358                    LevelCacheHint::L2CacheHint => {
359                        push_directive(tokens, "L2::cache_hint");
360                    }
361                }
362            }
363            if spaced {
364                tokens.push(PtxToken::Space);
365            }
366            self.tensormap.unparse_tokens_mode(tokens, spaced);
367            tokens.push(PtxToken::Comma);
368            if spaced {
369                tokens.push(PtxToken::Space);
370            }
371            self.srcmem.unparse_tokens_mode(tokens, spaced);
372            if self.cache_policy.is_some() {
373                tokens.push(PtxToken::Comma);
374            }
375            if let Some(opt_14) = self.cache_policy.as_ref() {
376                if spaced {
377                    tokens.push(PtxToken::Space);
378                }
379                opt_14.unparse_tokens_mode(tokens, spaced);
380            }
381            tokens.push(PtxToken::Semicolon);
382            if spaced {
383                tokens.push(PtxToken::Newline);
384            }
385        }
386    }
387}
ptx_parser/unparser/instruction/cp_async_bulk_tensor.rs

ptx_parser/unparser/instruction/
cp_async_bulk_tensor.rs
