ptx_parser/unparser/instruction/
cp_async_bulk.rs

1//! Original PTX specification:
2//!
3//! // global -> shared::cta
4//! cp.async.bulk.dst.src.completion_mechanism{.level::cache_hint} [dstMem], [srcMem], size, [mbar] {, cache-policy};
5//! .dst =                  { .shared::cta };
6//! .src =                  { .global };
7//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
8//! .level::cache_hint =    { .L2::cache_hint };
9//! ----------------------------------------------------------------
10//! // global -> shared::cluster;
11//! cp.async.bulk.dst.src.completion_mechanism{.multicast}{.level::cache_hint} [dstMem], [srcMem], size, [mbar] {, ctaMask} {, cache-policy};
12//! .dst =                  { .shared::cluster };
13//! .src =                  { .global };
14//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
15//! .level::cache_hint =    { .L2::cache_hint };
16//! .multicast =            { .multicast::cluster  };
17//! ----------------------------------------------------------------
18//! // shared::cta -> shared::cluster
19//! cp.async.bulk.dst.src.completion_mechanism [dstMem], [srcMem], size, [mbar];
20//! .dst =                  { .shared::cluster };
21//! .src =                  { .shared::cta };
22//! .completion_mechanism = { .mbarrier::complete_tx::bytes };
23//! ----------------------------------------------------------------
24//! // shared::cta -> global
25//! cp.async.bulk.dst.src.completion_mechanism{.level::cache_hint}{.cp_mask} [dstMem], [srcMem], size {, cache-policy} {, byteMask};
26//! .dst =                  { .global };
27//! .src =                  { .shared::cta };
28//! .completion_mechanism = { .bulk_group };
29//! .level::cache_hint =    { .L2::cache_hint };
30
31#![allow(unused)]
32
33use crate::lexer::PtxToken;
34use crate::unparser::{PtxUnparser, common::*};
35
36pub mod section_0 {
37    use super::*;
38    use crate::r#type::instruction::cp_async_bulk::section_0::*;
39
40    impl PtxUnparser for CpAsyncBulkDstSrcCompletionMechanismLevelCacheHint {
41        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
42            push_opcode(tokens, "cp");
43            push_directive(tokens, "async");
44            push_directive(tokens, "bulk");
45            match &self.dst {
46                Dst::SharedCta => {
47                    push_directive(tokens, "shared::cta");
48                }
49            }
50            match &self.src {
51                Src::Global => {
52                    push_directive(tokens, "global");
53                }
54            }
55            match &self.completion_mechanism {
56                CompletionMechanism::MbarrierCompleteTxBytes => {
57                    push_directive(tokens, "mbarrier::complete_tx::bytes");
58                }
59            }
60            if let Some(level_cache_hint_0) = self.level_cache_hint.as_ref() {
61                match level_cache_hint_0 {
62                    LevelCacheHint::L2CacheHint => {
63                        push_directive(tokens, "L2::cache_hint");
64                    }
65                }
66            }
67            self.dstmem.unparse_tokens(tokens);
68            tokens.push(PtxToken::Comma);
69            self.srcmem.unparse_tokens(tokens);
70            tokens.push(PtxToken::Comma);
71            self.size.unparse_tokens(tokens);
72            tokens.push(PtxToken::Comma);
73            self.mbar.unparse_tokens(tokens);
74            if self.cache_policy.is_some() {
75                tokens.push(PtxToken::Comma);
76            }
77            if let Some(opt_1) = self.cache_policy.as_ref() {
78                opt_1.unparse_tokens(tokens);
79            }
80            tokens.push(PtxToken::Semicolon);
81        }
82    }
83}
84
85pub mod section_1 {
86    use super::*;
87    use crate::r#type::instruction::cp_async_bulk::section_1::*;
88
89    impl PtxUnparser for CpAsyncBulkDstSrcCompletionMechanismMulticastLevelCacheHint {
90        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
91            push_opcode(tokens, "cp");
92            push_directive(tokens, "async");
93            push_directive(tokens, "bulk");
94            match &self.dst {
95                Dst::SharedCluster => {
96                    push_directive(tokens, "shared::cluster");
97                }
98            }
99            match &self.src {
100                Src::Global => {
101                    push_directive(tokens, "global");
102                }
103            }
104            match &self.completion_mechanism {
105                CompletionMechanism::MbarrierCompleteTxBytes => {
106                    push_directive(tokens, "mbarrier::complete_tx::bytes");
107                }
108            }
109            if let Some(multicast_2) = self.multicast.as_ref() {
110                match multicast_2 {
111                    Multicast::MulticastCluster => {
112                        push_directive(tokens, "multicast::cluster");
113                    }
114                }
115            }
116            if let Some(level_cache_hint_3) = self.level_cache_hint.as_ref() {
117                match level_cache_hint_3 {
118                    LevelCacheHint::L2CacheHint => {
119                        push_directive(tokens, "L2::cache_hint");
120                    }
121                }
122            }
123            self.dstmem.unparse_tokens(tokens);
124            tokens.push(PtxToken::Comma);
125            self.srcmem.unparse_tokens(tokens);
126            tokens.push(PtxToken::Comma);
127            self.size.unparse_tokens(tokens);
128            tokens.push(PtxToken::Comma);
129            self.mbar.unparse_tokens(tokens);
130            if self.ctamask.is_some() {
131                tokens.push(PtxToken::Comma);
132            }
133            if let Some(opt_4) = self.ctamask.as_ref() {
134                opt_4.unparse_tokens(tokens);
135            }
136            if self.cache_policy.is_some() {
137                tokens.push(PtxToken::Comma);
138            }
139            if let Some(opt_5) = self.cache_policy.as_ref() {
140                opt_5.unparse_tokens(tokens);
141            }
142            tokens.push(PtxToken::Semicolon);
143        }
144    }
145}
146
147pub mod section_2 {
148    use super::*;
149    use crate::r#type::instruction::cp_async_bulk::section_2::*;
150
151    impl PtxUnparser for CpAsyncBulkDstSrcCompletionMechanism {
152        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
153            push_opcode(tokens, "cp");
154            push_directive(tokens, "async");
155            push_directive(tokens, "bulk");
156            match &self.dst {
157                Dst::SharedCluster => {
158                    push_directive(tokens, "shared::cluster");
159                }
160            }
161            match &self.src {
162                Src::SharedCta => {
163                    push_directive(tokens, "shared::cta");
164                }
165            }
166            match &self.completion_mechanism {
167                CompletionMechanism::MbarrierCompleteTxBytes => {
168                    push_directive(tokens, "mbarrier::complete_tx::bytes");
169                }
170            }
171            self.dstmem.unparse_tokens(tokens);
172            tokens.push(PtxToken::Comma);
173            self.srcmem.unparse_tokens(tokens);
174            tokens.push(PtxToken::Comma);
175            self.size.unparse_tokens(tokens);
176            tokens.push(PtxToken::Comma);
177            self.mbar.unparse_tokens(tokens);
178            tokens.push(PtxToken::Semicolon);
179        }
180    }
181}
182
183pub mod section_3 {
184    use super::*;
185    use crate::r#type::instruction::cp_async_bulk::section_3::*;
186
187    impl PtxUnparser for CpAsyncBulkDstSrcCompletionMechanismLevelCacheHintCpMask {
188        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
189            push_opcode(tokens, "cp");
190            push_directive(tokens, "async");
191            push_directive(tokens, "bulk");
192            match &self.dst {
193                Dst::Global => {
194                    push_directive(tokens, "global");
195                }
196            }
197            match &self.src {
198                Src::SharedCta => {
199                    push_directive(tokens, "shared::cta");
200                }
201            }
202            match &self.completion_mechanism {
203                CompletionMechanism::BulkGroup => {
204                    push_directive(tokens, "bulk_group");
205                }
206            }
207            if let Some(level_cache_hint_6) = self.level_cache_hint.as_ref() {
208                match level_cache_hint_6 {
209                    LevelCacheHint::L2CacheHint => {
210                        push_directive(tokens, "L2::cache_hint");
211                    }
212                }
213            }
214            if self.cp_mask {
215                push_directive(tokens, "cp_mask");
216            }
217            self.dstmem.unparse_tokens(tokens);
218            tokens.push(PtxToken::Comma);
219            self.srcmem.unparse_tokens(tokens);
220            tokens.push(PtxToken::Comma);
221            self.size.unparse_tokens(tokens);
222            if self.cache_policy.is_some() {
223                tokens.push(PtxToken::Comma);
224            }
225            if let Some(opt_7) = self.cache_policy.as_ref() {
226                opt_7.unparse_tokens(tokens);
227            }
228            if self.bytemask.is_some() {
229                tokens.push(PtxToken::Comma);
230            }
231            if let Some(opt_8) = self.bytemask.as_ref() {
232                opt_8.unparse_tokens(tokens);
233            }
234            tokens.push(PtxToken::Semicolon);
235        }
236    }
237}
ptx_parser/unparser/instruction/cp_async_bulk.rs

ptx_parser/unparser/instruction/
cp_async_bulk.rs
