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            push_opcode(tokens, "cp");
45            push_directive(tokens, "async");
46            push_directive(tokens, "bulk");
47            push_directive(tokens, "tensor");
48            match &self.dim {
49                Dim::_1d => {
50                    push_directive(tokens, "1d");
51                }
52                Dim::_2d => {
53                    push_directive(tokens, "2d");
54                }
55                Dim::_3d => {
56                    push_directive(tokens, "3d");
57                }
58                Dim::_4d => {
59                    push_directive(tokens, "4d");
60                }
61                Dim::_5d => {
62                    push_directive(tokens, "5d");
63                }
64            }
65            match &self.dst {
66                Dst::SharedCta => {
67                    push_directive(tokens, "shared::cta");
68                }
69            }
70            match &self.src {
71                Src::Global => {
72                    push_directive(tokens, "global");
73                }
74            }
75            if let Some(load_mode_0) = self.load_mode.as_ref() {
76                match load_mode_0 {
77                    LoadMode::Im2colW128 => {
78                        push_directive(tokens, "im2col::w::128");
79                    }
80                    LoadMode::TileGather4 => {
81                        push_directive(tokens, "tile::gather4");
82                    }
83                    LoadMode::Im2colW => {
84                        push_directive(tokens, "im2col::w");
85                    }
86                    LoadMode::Im2col => {
87                        push_directive(tokens, "im2col");
88                    }
89                    LoadMode::Tile => {
90                        push_directive(tokens, "tile");
91                    }
92                }
93            }
94            match &self.completion_mechanism {
95                CompletionMechanism::MbarrierCompleteTxBytes => {
96                    push_directive(tokens, "mbarrier::complete_tx::bytes");
97                }
98            }
99            if let Some(cta_group_1) = self.cta_group.as_ref() {
100                match cta_group_1 {
101                    CtaGroup::CtaGroup1 => {
102                        push_directive(tokens, "cta_group::1");
103                    }
104                    CtaGroup::CtaGroup2 => {
105                        push_directive(tokens, "cta_group::2");
106                    }
107                }
108            }
109            if let Some(level_cache_hint_2) = self.level_cache_hint.as_ref() {
110                match level_cache_hint_2 {
111                    LevelCacheHint::L2CacheHint => {
112                        push_directive(tokens, "L2::cache_hint");
113                    }
114                }
115            }
116            self.dstmem.unparse_tokens(tokens);
117            tokens.push(PtxToken::Comma);
118            self.tensormap.unparse_tokens(tokens);
119            tokens.push(PtxToken::Comma);
120            self.mbar.unparse_tokens(tokens);
121            if self.im2colinfo.is_some() {
122                tokens.push(PtxToken::Comma);
123            }
124            if let Some(opt_3) = self.im2colinfo.as_ref() {
125                opt_3.unparse_tokens(tokens);
126            }
127            if self.cache_policy.is_some() {
128                tokens.push(PtxToken::Comma);
129            }
130            if let Some(opt_4) = self.cache_policy.as_ref() {
131                opt_4.unparse_tokens(tokens);
132            }
133            tokens.push(PtxToken::Semicolon);
134        }
135    }
136}
137
138pub mod section_1 {
139    use super::*;
140    use crate::r#type::instruction::cp_async_bulk_tensor::section_1::*;
141
142    impl PtxUnparser
143        for CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismMulticastCtaGroupLevelCacheHint
144    {
145        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
146            push_opcode(tokens, "cp");
147            push_directive(tokens, "async");
148            push_directive(tokens, "bulk");
149            push_directive(tokens, "tensor");
150            match &self.dim {
151                Dim::_1d => {
152                    push_directive(tokens, "1d");
153                }
154                Dim::_2d => {
155                    push_directive(tokens, "2d");
156                }
157                Dim::_3d => {
158                    push_directive(tokens, "3d");
159                }
160                Dim::_4d => {
161                    push_directive(tokens, "4d");
162                }
163                Dim::_5d => {
164                    push_directive(tokens, "5d");
165                }
166            }
167            match &self.dst {
168                Dst::SharedCluster => {
169                    push_directive(tokens, "shared::cluster");
170                }
171            }
172            match &self.src {
173                Src::Global => {
174                    push_directive(tokens, "global");
175                }
176            }
177            if let Some(load_mode_5) = self.load_mode.as_ref() {
178                match load_mode_5 {
179                    LoadMode::Im2colW128 => {
180                        push_directive(tokens, "im2col::w::128");
181                    }
182                    LoadMode::TileGather4 => {
183                        push_directive(tokens, "tile::gather4");
184                    }
185                    LoadMode::Im2colW => {
186                        push_directive(tokens, "im2col::w");
187                    }
188                    LoadMode::Im2col => {
189                        push_directive(tokens, "im2col");
190                    }
191                    LoadMode::Tile => {
192                        push_directive(tokens, "tile");
193                    }
194                }
195            }
196            match &self.completion_mechanism {
197                CompletionMechanism::MbarrierCompleteTxBytes => {
198                    push_directive(tokens, "mbarrier::complete_tx::bytes");
199                }
200            }
201            if let Some(multicast_6) = self.multicast.as_ref() {
202                match multicast_6 {
203                    Multicast::MulticastCluster => {
204                        push_directive(tokens, "multicast::cluster");
205                    }
206                }
207            }
208            if let Some(cta_group_7) = self.cta_group.as_ref() {
209                match cta_group_7 {
210                    CtaGroup::CtaGroup1 => {
211                        push_directive(tokens, "cta_group::1");
212                    }
213                    CtaGroup::CtaGroup2 => {
214                        push_directive(tokens, "cta_group::2");
215                    }
216                }
217            }
218            if let Some(level_cache_hint_8) = self.level_cache_hint.as_ref() {
219                match level_cache_hint_8 {
220                    LevelCacheHint::L2CacheHint => {
221                        push_directive(tokens, "L2::cache_hint");
222                    }
223                }
224            }
225            self.dstmem.unparse_tokens(tokens);
226            tokens.push(PtxToken::Comma);
227            self.tensormap.unparse_tokens(tokens);
228            tokens.push(PtxToken::Comma);
229            self.mbar.unparse_tokens(tokens);
230            if self.im2colinfo.is_some() {
231                tokens.push(PtxToken::Comma);
232            }
233            if let Some(opt_9) = self.im2colinfo.as_ref() {
234                opt_9.unparse_tokens(tokens);
235            }
236            if self.ctamask.is_some() {
237                tokens.push(PtxToken::Comma);
238            }
239            if let Some(opt_10) = self.ctamask.as_ref() {
240                opt_10.unparse_tokens(tokens);
241            }
242            if self.cache_policy.is_some() {
243                tokens.push(PtxToken::Comma);
244            }
245            if let Some(opt_11) = self.cache_policy.as_ref() {
246                opt_11.unparse_tokens(tokens);
247            }
248            tokens.push(PtxToken::Semicolon);
249        }
250    }
251}
252
253pub mod section_2 {
254    use super::*;
255    use crate::r#type::instruction::cp_async_bulk_tensor::section_2::*;
256
257    impl PtxUnparser for CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismLevelCacheHint {
258        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
259            push_opcode(tokens, "cp");
260            push_directive(tokens, "async");
261            push_directive(tokens, "bulk");
262            push_directive(tokens, "tensor");
263            match &self.dim {
264                Dim::_1d => {
265                    push_directive(tokens, "1d");
266                }
267                Dim::_2d => {
268                    push_directive(tokens, "2d");
269                }
270                Dim::_3d => {
271                    push_directive(tokens, "3d");
272                }
273                Dim::_4d => {
274                    push_directive(tokens, "4d");
275                }
276                Dim::_5d => {
277                    push_directive(tokens, "5d");
278                }
279            }
280            match &self.dst {
281                Dst::Global => {
282                    push_directive(tokens, "global");
283                }
284            }
285            match &self.src {
286                Src::SharedCta => {
287                    push_directive(tokens, "shared::cta");
288                }
289            }
290            if let Some(load_mode_12) = self.load_mode.as_ref() {
291                match load_mode_12 {
292                    LoadMode::TileScatter4 => {
293                        push_directive(tokens, "tile::scatter4");
294                    }
295                    LoadMode::Im2colNoOffs => {
296                        push_directive(tokens, "im2col_no_offs");
297                    }
298                    LoadMode::Tile => {
299                        push_directive(tokens, "tile");
300                    }
301                }
302            }
303            match &self.completion_mechanism {
304                CompletionMechanism::BulkGroup => {
305                    push_directive(tokens, "bulk_group");
306                }
307            }
308            if let Some(level_cache_hint_13) = self.level_cache_hint.as_ref() {
309                match level_cache_hint_13 {
310                    LevelCacheHint::L2CacheHint => {
311                        push_directive(tokens, "L2::cache_hint");
312                    }
313                }
314            }
315            self.tensormap.unparse_tokens(tokens);
316            tokens.push(PtxToken::Comma);
317            self.srcmem.unparse_tokens(tokens);
318            if self.cache_policy.is_some() {
319                tokens.push(PtxToken::Comma);
320            }
321            if let Some(opt_14) = self.cache_policy.as_ref() {
322                opt_14.unparse_tokens(tokens);
323            }
324            tokens.push(PtxToken::Semicolon);
325        }
326    }
327}
ptx_parser/unparser/instruction/cp_async_bulk_tensor.rs

ptx_parser/unparser/instruction/
cp_async_bulk_tensor.rs
