ptx_parser/parser/instruction/
tcgen05_shift.rs1#![allow(unused)]
7
8use crate::parser::{
9 PtxParseError, PtxParser, PtxTokenStream, Span,
10 util::{
11 between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
12 pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
13 },
14};
15use crate::r#type::common::*;
16use crate::{alt, ok, seq_n};
17
18pub mod section_0 {
19 use super::*;
20 use crate::r#type::instruction::tcgen05_shift::section_0::*;
21
22 impl PtxParser for CtaGroup {
27 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
28 alt!(
29 map(string_p(".cta_group::1"), |_, _span| CtaGroup::CtaGroup1),
30 map(string_p(".cta_group::2"), |_, _span| CtaGroup::CtaGroup2)
31 )
32 }
33 }
34
35 impl PtxParser for Tcgen05ShiftCtaGroupDown {
36 fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
37 try_map(
38 seq_n!(
39 string_p("tcgen05"),
40 string_p(".shift"),
41 CtaGroup::parse(),
42 string_p(".down"),
43 AddressOperand::parse(),
44 semicolon_p()
45 ),
46 |(_, shift, cta_group, down, taddr, _), span| {
47 ok!(Tcgen05ShiftCtaGroupDown {
48 shift = shift,
49 cta_group = cta_group,
50 down = down,
51 taddr = taddr,
52
53 })
54 },
55 )
56 }
57 }
58}