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() { tokens.push(PtxToken::Comma); }
75                    if let Some(opt_1) = self.cache_policy.as_ref() {
76                        opt_1.unparse_tokens(tokens);
77                    }
78            tokens.push(PtxToken::Semicolon);
79        }
80    }
81
82}
83
84pub mod section_1 {
85    use super::*;
86    use crate::r#type::instruction::cp_async_bulk::section_1::*;
87
88    impl PtxUnparser for CpAsyncBulkDstSrcCompletionMechanismMulticastLevelCacheHint {
89        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
90            push_opcode(tokens, "cp");
91                    push_directive(tokens, "async");
92                    push_directive(tokens, "bulk");
93                    match &self.dst {
94                            Dst::SharedCluster => {
95                                    push_directive(tokens, "shared::cluster");
96                            }
97                    }
98                    match &self.src {
99                            Src::Global => {
100                                    push_directive(tokens, "global");
101                            }
102                    }
103                    match &self.completion_mechanism {
104                            CompletionMechanism::MbarrierCompleteTxBytes => {
105                                    push_directive(tokens, "mbarrier::complete_tx::bytes");
106                            }
107                    }
108                    if let Some(multicast_2) = self.multicast.as_ref() {
109                            match multicast_2 {
110                                    Multicast::MulticastCluster => {
111                                            push_directive(tokens, "multicast::cluster");
112                                    }
113                            }
114                    }
115                    if let Some(level_cache_hint_3) = self.level_cache_hint.as_ref() {
116                            match level_cache_hint_3 {
117                                    LevelCacheHint::L2CacheHint => {
118                                            push_directive(tokens, "L2::cache_hint");
119                                    }
120                            }
121                    }
122                    self.dstmem.unparse_tokens(tokens);
123            tokens.push(PtxToken::Comma);
124                    self.srcmem.unparse_tokens(tokens);
125            tokens.push(PtxToken::Comma);
126                    self.size.unparse_tokens(tokens);
127            tokens.push(PtxToken::Comma);
128                    self.mbar.unparse_tokens(tokens);
129            if self.ctamask.is_some() { tokens.push(PtxToken::Comma); }
130                    if let Some(opt_4) = self.ctamask.as_ref() {
131                        opt_4.unparse_tokens(tokens);
132                    }
133            if self.cache_policy.is_some() { tokens.push(PtxToken::Comma); }
134                    if let Some(opt_5) = self.cache_policy.as_ref() {
135                        opt_5.unparse_tokens(tokens);
136                    }
137            tokens.push(PtxToken::Semicolon);
138        }
139    }
140
141}
142
143pub mod section_2 {
144    use super::*;
145    use crate::r#type::instruction::cp_async_bulk::section_2::*;
146
147    impl PtxUnparser for CpAsyncBulkDstSrcCompletionMechanism {
148        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
149            push_opcode(tokens, "cp");
150                    push_directive(tokens, "async");
151                    push_directive(tokens, "bulk");
152                    match &self.dst {
153                            Dst::SharedCluster => {
154                                    push_directive(tokens, "shared::cluster");
155                            }
156                    }
157                    match &self.src {
158                            Src::SharedCta => {
159                                    push_directive(tokens, "shared::cta");
160                            }
161                    }
162                    match &self.completion_mechanism {
163                            CompletionMechanism::MbarrierCompleteTxBytes => {
164                                    push_directive(tokens, "mbarrier::complete_tx::bytes");
165                            }
166                    }
167                    self.dstmem.unparse_tokens(tokens);
168            tokens.push(PtxToken::Comma);
169                    self.srcmem.unparse_tokens(tokens);
170            tokens.push(PtxToken::Comma);
171                    self.size.unparse_tokens(tokens);
172            tokens.push(PtxToken::Comma);
173                    self.mbar.unparse_tokens(tokens);
174            tokens.push(PtxToken::Semicolon);
175        }
176    }
177
178}
179
180pub mod section_3 {
181    use super::*;
182    use crate::r#type::instruction::cp_async_bulk::section_3::*;
183
184    impl PtxUnparser for CpAsyncBulkDstSrcCompletionMechanismLevelCacheHintCpMask {
185        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
186            push_opcode(tokens, "cp");
187                    push_directive(tokens, "async");
188                    push_directive(tokens, "bulk");
189                    match &self.dst {
190                            Dst::Global => {
191                                    push_directive(tokens, "global");
192                            }
193                    }
194                    match &self.src {
195                            Src::SharedCta => {
196                                    push_directive(tokens, "shared::cta");
197                            }
198                    }
199                    match &self.completion_mechanism {
200                            CompletionMechanism::BulkGroup => {
201                                    push_directive(tokens, "bulk_group");
202                            }
203                    }
204                    if let Some(level_cache_hint_6) = self.level_cache_hint.as_ref() {
205                            match level_cache_hint_6 {
206                                    LevelCacheHint::L2CacheHint => {
207                                            push_directive(tokens, "L2::cache_hint");
208                                    }
209                            }
210                    }
211                    if self.cp_mask {
212                            push_directive(tokens, "cp_mask");
213                    }
214                    self.dstmem.unparse_tokens(tokens);
215            tokens.push(PtxToken::Comma);
216                    self.srcmem.unparse_tokens(tokens);
217            tokens.push(PtxToken::Comma);
218                    self.size.unparse_tokens(tokens);
219            if self.cache_policy.is_some() { tokens.push(PtxToken::Comma); }
220                    if let Some(opt_7) = self.cache_policy.as_ref() {
221                        opt_7.unparse_tokens(tokens);
222                    }
223            if self.bytemask.is_some() { tokens.push(PtxToken::Comma); }
224                    if let Some(opt_8) = self.bytemask.as_ref() {
225                        opt_8.unparse_tokens(tokens);
226                    }
227            tokens.push(PtxToken::Semicolon);
228        }
229    }
230
231}
232
ptx_parser/unparser/instruction/cp_async_bulk.rs

ptx_parser/unparser/instruction/
cp_async_bulk.rs
