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() { tokens.push(PtxToken::Comma); }
122 if let Some(opt_3) = self.im2colinfo.as_ref() {
123 opt_3.unparse_tokens(tokens);
124 }
125 if self.cache_policy.is_some() { tokens.push(PtxToken::Comma); }
126 if let Some(opt_4) = self.cache_policy.as_ref() {
127 opt_4.unparse_tokens(tokens);
128 }
129 tokens.push(PtxToken::Semicolon);
130 }
131 }
132
133}
134
135pub mod section_1 {
136 use super::*;
137 use crate::r#type::instruction::cp_async_bulk_tensor::section_1::*;
138
139 impl PtxUnparser for CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismMulticastCtaGroupLevelCacheHint {
140 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
141 push_opcode(tokens, "cp");
142 push_directive(tokens, "async");
143 push_directive(tokens, "bulk");
144 push_directive(tokens, "tensor");
145 match &self.dim {
146 Dim::_1d => {
147 push_directive(tokens, "1d");
148 }
149 Dim::_2d => {
150 push_directive(tokens, "2d");
151 }
152 Dim::_3d => {
153 push_directive(tokens, "3d");
154 }
155 Dim::_4d => {
156 push_directive(tokens, "4d");
157 }
158 Dim::_5d => {
159 push_directive(tokens, "5d");
160 }
161 }
162 match &self.dst {
163 Dst::SharedCluster => {
164 push_directive(tokens, "shared::cluster");
165 }
166 }
167 match &self.src {
168 Src::Global => {
169 push_directive(tokens, "global");
170 }
171 }
172 if let Some(load_mode_5) = self.load_mode.as_ref() {
173 match load_mode_5 {
174 LoadMode::Im2colW128 => {
175 push_directive(tokens, "im2col::w::128");
176 }
177 LoadMode::TileGather4 => {
178 push_directive(tokens, "tile::gather4");
179 }
180 LoadMode::Im2colW => {
181 push_directive(tokens, "im2col::w");
182 }
183 LoadMode::Im2col => {
184 push_directive(tokens, "im2col");
185 }
186 LoadMode::Tile => {
187 push_directive(tokens, "tile");
188 }
189 }
190 }
191 match &self.completion_mechanism {
192 CompletionMechanism::MbarrierCompleteTxBytes => {
193 push_directive(tokens, "mbarrier::complete_tx::bytes");
194 }
195 }
196 if let Some(multicast_6) = self.multicast.as_ref() {
197 match multicast_6 {
198 Multicast::MulticastCluster => {
199 push_directive(tokens, "multicast::cluster");
200 }
201 }
202 }
203 if let Some(cta_group_7) = self.cta_group.as_ref() {
204 match cta_group_7 {
205 CtaGroup::CtaGroup1 => {
206 push_directive(tokens, "cta_group::1");
207 }
208 CtaGroup::CtaGroup2 => {
209 push_directive(tokens, "cta_group::2");
210 }
211 }
212 }
213 if let Some(level_cache_hint_8) = self.level_cache_hint.as_ref() {
214 match level_cache_hint_8 {
215 LevelCacheHint::L2CacheHint => {
216 push_directive(tokens, "L2::cache_hint");
217 }
218 }
219 }
220 self.dstmem.unparse_tokens(tokens);
221 tokens.push(PtxToken::Comma);
222 self.tensormap.unparse_tokens(tokens);
223 tokens.push(PtxToken::Comma);
224 self.mbar.unparse_tokens(tokens);
225 if self.im2colinfo.is_some() { tokens.push(PtxToken::Comma); }
226 if let Some(opt_9) = self.im2colinfo.as_ref() {
227 opt_9.unparse_tokens(tokens);
228 }
229 if self.ctamask.is_some() { tokens.push(PtxToken::Comma); }
230 if let Some(opt_10) = self.ctamask.as_ref() {
231 opt_10.unparse_tokens(tokens);
232 }
233 if self.cache_policy.is_some() { tokens.push(PtxToken::Comma); }
234 if let Some(opt_11) = self.cache_policy.as_ref() {
235 opt_11.unparse_tokens(tokens);
236 }
237 tokens.push(PtxToken::Semicolon);
238 }
239 }
240
241}
242
243pub mod section_2 {
244 use super::*;
245 use crate::r#type::instruction::cp_async_bulk_tensor::section_2::*;
246
247 impl PtxUnparser for CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismLevelCacheHint {
248 fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
249 push_opcode(tokens, "cp");
250 push_directive(tokens, "async");
251 push_directive(tokens, "bulk");
252 push_directive(tokens, "tensor");
253 match &self.dim {
254 Dim::_1d => {
255 push_directive(tokens, "1d");
256 }
257 Dim::_2d => {
258 push_directive(tokens, "2d");
259 }
260 Dim::_3d => {
261 push_directive(tokens, "3d");
262 }
263 Dim::_4d => {
264 push_directive(tokens, "4d");
265 }
266 Dim::_5d => {
267 push_directive(tokens, "5d");
268 }
269 }
270 match &self.dst {
271 Dst::Global => {
272 push_directive(tokens, "global");
273 }
274 }
275 match &self.src {
276 Src::SharedCta => {
277 push_directive(tokens, "shared::cta");
278 }
279 }
280 if let Some(load_mode_12) = self.load_mode.as_ref() {
281 match load_mode_12 {
282 LoadMode::TileScatter4 => {
283 push_directive(tokens, "tile::scatter4");
284 }
285 LoadMode::Im2colNoOffs => {
286 push_directive(tokens, "im2col_no_offs");
287 }
288 LoadMode::Tile => {
289 push_directive(tokens, "tile");
290 }
291 }
292 }
293 match &self.completion_mechanism {
294 CompletionMechanism::BulkGroup => {
295 push_directive(tokens, "bulk_group");
296 }
297 }
298 if let Some(level_cache_hint_13) = self.level_cache_hint.as_ref() {
299 match level_cache_hint_13 {
300 LevelCacheHint::L2CacheHint => {
301 push_directive(tokens, "L2::cache_hint");
302 }
303 }
304 }
305 self.tensormap.unparse_tokens(tokens);
306 tokens.push(PtxToken::Comma);
307 self.srcmem.unparse_tokens(tokens);
308 if self.cache_policy.is_some() { tokens.push(PtxToken::Comma); }
309 if let Some(opt_14) = self.cache_policy.as_ref() {
310 opt_14.unparse_tokens(tokens);
311 }
312 tokens.push(PtxToken::Semicolon);
313 }
314 }
315
316}
317