ptx_parser/parser/instruction/
tcgen05_alloc.rs

1//! Original PTX specification:
2//!
3//! tcgen05.alloc.cta_group.sync.aligned{.shared::cta}.b32  [dst], nCols;
4//! tcgen05.dealloc.cta_group.sync.aligned.b32              taddr, nCols;
5//! tcgen05.relinquish_alloc_permit.cta_group.sync.aligned;
6//! .cta_group = { .cta_group::1, .cta_group::2 };
7
8#![allow(unused)]
9
10use crate::parser::{
11    PtxParseError, PtxParser, PtxTokenStream, Span,
12    util::{
13        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
14        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
15    },
16};
17use crate::r#type::common::*;
18use crate::{alt, ok, seq_n};
19
20pub mod section_0 {
21    use super::*;
22    use crate::r#type::instruction::tcgen05_alloc::section_0::*;
23
24    // ============================================================================
25    // Generated enum parsers
26    // ============================================================================
27
28    impl PtxParser for CtaGroup {
29        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
30            alt!(
31                map(string_p(".cta_group::1"), |_, _span| CtaGroup::CtaGroup1),
32                map(string_p(".cta_group::2"), |_, _span| CtaGroup::CtaGroup2)
33            )
34        }
35    }
36
37    impl PtxParser for Tcgen05AllocCtaGroupSyncAlignedSharedCtaB32 {
38        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
39            try_map(
40                seq_n!(
41                    string_p("tcgen05"),
42                    string_p(".alloc"),
43                    CtaGroup::parse(),
44                    string_p(".sync"),
45                    string_p(".aligned"),
46                    map(optional(string_p(".shared::cta")), |value, _| value
47                        .is_some()),
48                    string_p(".b32"),
49                    AddressOperand::parse(),
50                    comma_p(),
51                    GeneralOperand::parse(),
52                    semicolon_p()
53                ),
54                |(_, alloc, cta_group, sync, aligned, shared_cta, b32, dst, _, ncols, _), span| {
55                    ok!(Tcgen05AllocCtaGroupSyncAlignedSharedCtaB32 {
56                        alloc = alloc,
57                        cta_group = cta_group,
58                        sync = sync,
59                        aligned = aligned,
60                        shared_cta = shared_cta,
61                        b32 = b32,
62                        dst = dst,
63                        ncols = ncols,
64
65                    })
66                },
67            )
68        }
69    }
70
71    impl PtxParser for Tcgen05DeallocCtaGroupSyncAlignedB32 {
72        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
73            try_map(
74                seq_n!(
75                    string_p("tcgen05"),
76                    string_p(".dealloc"),
77                    CtaGroup::parse(),
78                    string_p(".sync"),
79                    string_p(".aligned"),
80                    string_p(".b32"),
81                    GeneralOperand::parse(),
82                    comma_p(),
83                    GeneralOperand::parse(),
84                    semicolon_p()
85                ),
86                |(_, dealloc, cta_group, sync, aligned, b32, taddr, _, ncols, _), span| {
87                    ok!(Tcgen05DeallocCtaGroupSyncAlignedB32 {
88                        dealloc = dealloc,
89                        cta_group = cta_group,
90                        sync = sync,
91                        aligned = aligned,
92                        b32 = b32,
93                        taddr = taddr,
94                        ncols = ncols,
95
96                    })
97                },
98            )
99        }
100    }
101
102    impl PtxParser for Tcgen05RelinquishAllocPermitCtaGroupSyncAligned {
103        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
104            try_map(
105                seq_n!(
106                    string_p("tcgen05"),
107                    string_p(".relinquish_alloc_permit"),
108                    CtaGroup::parse(),
109                    string_p(".sync"),
110                    string_p(".aligned"),
111                    semicolon_p()
112                ),
113                |(_, relinquish_alloc_permit, cta_group, sync, aligned, _), span| {
114                    ok!(Tcgen05RelinquishAllocPermitCtaGroupSyncAligned {
115                        relinquish_alloc_permit = relinquish_alloc_permit,
116                        cta_group = cta_group,
117                        sync = sync,
118                        aligned = aligned,
119
120                    })
121                },
122            )
123        }
124    }
125}
ptx_parser/parser/instruction/tcgen05_alloc.rs

ptx_parser/parser/instruction/
tcgen05_alloc.rs
