Skip to main content

ptx_parser/unparser/instruction/
tcgen05_shift.rs

1//! Original PTX specification:
2//!
3//! tcgen05.shift.cta_group.down  [taddr];
4//! .cta_group = { .cta_group::1, .cta_group::2 }
5
6#![allow(unused)]
7
8use crate::lexer::PtxToken;
9use crate::unparser::{PtxUnparser, common::*};
10
11pub mod section_0 {
12    use super::*;
13    use crate::r#type::instruction::tcgen05_shift::section_0::*;
14
15    impl PtxUnparser for Tcgen05ShiftCtaGroupDown {
16        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
17            self.unparse_tokens_mode(tokens, false);
18        }
19        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
20            push_opcode(tokens, "tcgen05");
21            push_directive(tokens, "shift");
22            match &self.cta_group {
23                CtaGroup::CtaGroup1 => {
24                    push_directive(tokens, "cta_group::1");
25                }
26                CtaGroup::CtaGroup2 => {
27                    push_directive(tokens, "cta_group::2");
28                }
29            }
30            push_directive(tokens, "down");
31            if spaced {
32                tokens.push(PtxToken::Space);
33            }
34            self.taddr.unparse_tokens_mode(tokens, spaced);
35            tokens.push(PtxToken::Semicolon);
36            if spaced {
37                tokens.push(PtxToken::Newline);
38            }
39        }
40    }
41}