ptx_parser/parser/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::parser::{
34    PtxParseError, PtxParser, PtxTokenStream, Span,
35    util::{
36        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
37        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
38    },
39};
40use crate::r#type::common::*;
41use crate::{alt, ok, seq_n};
42
43pub mod section_0 {
44    use super::*;
45    use crate::r#type::instruction::cp_async_bulk::section_0::*;
46
47    // ============================================================================
48    // Generated enum parsers
49    // ============================================================================
50
51    impl PtxParser for CompletionMechanism {
52        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
53            alt!(map(
54                string_p(".mbarrier::complete_tx::bytes"),
55                |_, _span| CompletionMechanism::MbarrierCompleteTxBytes
56            ))
57        }
58    }
59
60    impl PtxParser for Dst {
61        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
62            alt!(map(string_p(".shared::cta"), |_, _span| Dst::SharedCta))
63        }
64    }
65
66    impl PtxParser for LevelCacheHint {
67        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
68            alt!(map(string_p(".L2::cache_hint"), |_, _span| {
69                LevelCacheHint::L2CacheHint
70            }))
71        }
72    }
73
74    impl PtxParser for Src {
75        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
76            alt!(map(string_p(".global"), |_, _span| Src::Global))
77        }
78    }
79
80    impl PtxParser for CpAsyncBulkDstSrcCompletionMechanismLevelCacheHint {
81        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
82            try_map(
83                seq_n!(
84                    string_p("cp"),
85                    string_p(".async"),
86                    string_p(".bulk"),
87                    Dst::parse(),
88                    Src::parse(),
89                    CompletionMechanism::parse(),
90                    optional(LevelCacheHint::parse()),
91                    AddressOperand::parse(),
92                    comma_p(),
93                    AddressOperand::parse(),
94                    comma_p(),
95                    GeneralOperand::parse(),
96                    comma_p(),
97                    AddressOperand::parse(),
98                    map(
99                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
100                        |value, _| value.map(|(_, operand)| operand)
101                    ),
102                    semicolon_p()
103                ),
104                |(
105                    _,
106                    async_,
107                    bulk,
108                    dst,
109                    src,
110                    completion_mechanism,
111                    level_cache_hint,
112                    dstmem,
113                    _,
114                    srcmem,
115                    _,
116                    size,
117                    _,
118                    mbar,
119                    cache_policy,
120                    _,
121                ),
122                 span| {
123                    ok!(CpAsyncBulkDstSrcCompletionMechanismLevelCacheHint {
124                        async_ = async_,
125                        bulk = bulk,
126                        dst = dst,
127                        src = src,
128                        completion_mechanism = completion_mechanism,
129                        level_cache_hint = level_cache_hint,
130                        dstmem = dstmem,
131                        srcmem = srcmem,
132                        size = size,
133                        mbar = mbar,
134                        cache_policy = cache_policy,
135
136                    })
137                },
138            )
139        }
140    }
141}
142
143pub mod section_1 {
144    use super::*;
145    use crate::r#type::instruction::cp_async_bulk::section_1::*;
146
147    // ============================================================================
148    // Generated enum parsers
149    // ============================================================================
150
151    impl PtxParser for CompletionMechanism {
152        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
153            alt!(map(
154                string_p(".mbarrier::complete_tx::bytes"),
155                |_, _span| CompletionMechanism::MbarrierCompleteTxBytes
156            ))
157        }
158    }
159
160    impl PtxParser for Dst {
161        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
162            alt!(map(string_p(".shared::cluster"), |_, _span| {
163                Dst::SharedCluster
164            }))
165        }
166    }
167
168    impl PtxParser for LevelCacheHint {
169        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
170            alt!(map(string_p(".L2::cache_hint"), |_, _span| {
171                LevelCacheHint::L2CacheHint
172            }))
173        }
174    }
175
176    impl PtxParser for Multicast {
177        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
178            alt!(map(string_p(".multicast::cluster"), |_, _span| {
179                Multicast::MulticastCluster
180            }))
181        }
182    }
183
184    impl PtxParser for Src {
185        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
186            alt!(map(string_p(".global"), |_, _span| Src::Global))
187        }
188    }
189
190    impl PtxParser for CpAsyncBulkDstSrcCompletionMechanismMulticastLevelCacheHint {
191        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
192            try_map(
193                seq_n!(
194                    string_p("cp"),
195                    string_p(".async"),
196                    string_p(".bulk"),
197                    Dst::parse(),
198                    Src::parse(),
199                    CompletionMechanism::parse(),
200                    optional(Multicast::parse()),
201                    optional(LevelCacheHint::parse()),
202                    AddressOperand::parse(),
203                    comma_p(),
204                    AddressOperand::parse(),
205                    comma_p(),
206                    GeneralOperand::parse(),
207                    comma_p(),
208                    AddressOperand::parse(),
209                    map(
210                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
211                        |value, _| value.map(|(_, operand)| operand)
212                    ),
213                    map(
214                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
215                        |value, _| value.map(|(_, operand)| operand)
216                    ),
217                    semicolon_p()
218                ),
219                |(
220                    _,
221                    async_,
222                    bulk,
223                    dst,
224                    src,
225                    completion_mechanism,
226                    multicast,
227                    level_cache_hint,
228                    dstmem,
229                    _,
230                    srcmem,
231                    _,
232                    size,
233                    _,
234                    mbar,
235                    ctamask,
236                    cache_policy,
237                    _,
238                ),
239                 span| {
240                    ok!(CpAsyncBulkDstSrcCompletionMechanismMulticastLevelCacheHint {
241                        async_ = async_,
242                        bulk = bulk,
243                        dst = dst,
244                        src = src,
245                        completion_mechanism = completion_mechanism,
246                        multicast = multicast,
247                        level_cache_hint = level_cache_hint,
248                        dstmem = dstmem,
249                        srcmem = srcmem,
250                        size = size,
251                        mbar = mbar,
252                        ctamask = ctamask,
253                        cache_policy = cache_policy,
254
255                    })
256                },
257            )
258        }
259    }
260}
261
262pub mod section_2 {
263    use super::*;
264    use crate::r#type::instruction::cp_async_bulk::section_2::*;
265
266    // ============================================================================
267    // Generated enum parsers
268    // ============================================================================
269
270    impl PtxParser for CompletionMechanism {
271        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
272            alt!(map(
273                string_p(".mbarrier::complete_tx::bytes"),
274                |_, _span| CompletionMechanism::MbarrierCompleteTxBytes
275            ))
276        }
277    }
278
279    impl PtxParser for Dst {
280        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
281            alt!(map(string_p(".shared::cluster"), |_, _span| {
282                Dst::SharedCluster
283            }))
284        }
285    }
286
287    impl PtxParser for Src {
288        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
289            alt!(map(string_p(".shared::cta"), |_, _span| Src::SharedCta))
290        }
291    }
292
293    impl PtxParser for CpAsyncBulkDstSrcCompletionMechanism {
294        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
295            try_map(
296                seq_n!(
297                    string_p("cp"),
298                    string_p(".async"),
299                    string_p(".bulk"),
300                    Dst::parse(),
301                    Src::parse(),
302                    CompletionMechanism::parse(),
303                    AddressOperand::parse(),
304                    comma_p(),
305                    AddressOperand::parse(),
306                    comma_p(),
307                    GeneralOperand::parse(),
308                    comma_p(),
309                    AddressOperand::parse(),
310                    semicolon_p()
311                ),
312                |(
313                    _,
314                    async_,
315                    bulk,
316                    dst,
317                    src,
318                    completion_mechanism,
319                    dstmem,
320                    _,
321                    srcmem,
322                    _,
323                    size,
324                    _,
325                    mbar,
326                    _,
327                ),
328                 span| {
329                    ok!(CpAsyncBulkDstSrcCompletionMechanism {
330                        async_ = async_,
331                        bulk = bulk,
332                        dst = dst,
333                        src = src,
334                        completion_mechanism = completion_mechanism,
335                        dstmem = dstmem,
336                        srcmem = srcmem,
337                        size = size,
338                        mbar = mbar,
339
340                    })
341                },
342            )
343        }
344    }
345}
346
347pub mod section_3 {
348    use super::*;
349    use crate::r#type::instruction::cp_async_bulk::section_3::*;
350
351    // ============================================================================
352    // Generated enum parsers
353    // ============================================================================
354
355    impl PtxParser for CompletionMechanism {
356        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
357            alt!(map(string_p(".bulk_group"), |_, _span| {
358                CompletionMechanism::BulkGroup
359            }))
360        }
361    }
362
363    impl PtxParser for Dst {
364        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
365            alt!(map(string_p(".global"), |_, _span| Dst::Global))
366        }
367    }
368
369    impl PtxParser for LevelCacheHint {
370        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
371            alt!(map(string_p(".L2::cache_hint"), |_, _span| {
372                LevelCacheHint::L2CacheHint
373            }))
374        }
375    }
376
377    impl PtxParser for Src {
378        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
379            alt!(map(string_p(".shared::cta"), |_, _span| Src::SharedCta))
380        }
381    }
382
383    impl PtxParser for CpAsyncBulkDstSrcCompletionMechanismLevelCacheHintCpMask {
384        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
385            try_map(
386                seq_n!(
387                    string_p("cp"),
388                    string_p(".async"),
389                    string_p(".bulk"),
390                    Dst::parse(),
391                    Src::parse(),
392                    CompletionMechanism::parse(),
393                    optional(LevelCacheHint::parse()),
394                    map(optional(string_p(".cp_mask")), |value, _| value.is_some()),
395                    AddressOperand::parse(),
396                    comma_p(),
397                    AddressOperand::parse(),
398                    comma_p(),
399                    GeneralOperand::parse(),
400                    map(
401                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
402                        |value, _| value.map(|(_, operand)| operand)
403                    ),
404                    map(
405                        optional(seq_n!(comma_p(), GeneralOperand::parse())),
406                        |value, _| value.map(|(_, operand)| operand)
407                    ),
408                    semicolon_p()
409                ),
410                |(
411                    _,
412                    async_,
413                    bulk,
414                    dst,
415                    src,
416                    completion_mechanism,
417                    level_cache_hint,
418                    cp_mask,
419                    dstmem,
420                    _,
421                    srcmem,
422                    _,
423                    size,
424                    cache_policy,
425                    bytemask,
426                    _,
427                ),
428                 span| {
429                    ok!(CpAsyncBulkDstSrcCompletionMechanismLevelCacheHintCpMask {
430                        async_ = async_,
431                        bulk = bulk,
432                        dst = dst,
433                        src = src,
434                        completion_mechanism = completion_mechanism,
435                        level_cache_hint = level_cache_hint,
436                        cp_mask = cp_mask,
437                        dstmem = dstmem,
438                        srcmem = srcmem,
439                        size = size,
440                        cache_policy = cache_policy,
441                        bytemask = bytemask,
442
443                    })
444                },
445            )
446        }
447    }
448}
ptx_parser/parser/instruction/cp_async_bulk.rs

ptx_parser/parser/instruction/
cp_async_bulk.rs
