ptx_parser/unparser/instruction/
stmatrix.rs

1//! Original PTX specification:
2//!
3//! stmatrix.sync.aligned.shape.num{.trans}{.ss}.type [p], r;
4//! .shape  = {.m8n8, .m16n8};
5//! .num    = {.x1, .x2, .x4};
6//! .ss     = {.shared, .shared::cta};
7//! .type   = {.b16, .b8};
8
9#![allow(unused)]
10
11use crate::lexer::PtxToken;
12use crate::unparser::{PtxUnparser, common::*};
13
14pub mod section_0 {
15    use super::*;
16    use crate::r#type::instruction::stmatrix::section_0::*;
17
18    impl PtxUnparser for StmatrixSyncAlignedShapeNumTransSsType {
19        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
20            push_opcode(tokens, "stmatrix");
21            push_directive(tokens, "sync");
22            push_directive(tokens, "aligned");
23            match &self.shape {
24                Shape::M16n8 => {
25                    push_directive(tokens, "m16n8");
26                }
27                Shape::M8n8 => {
28                    push_directive(tokens, "m8n8");
29                }
30            }
31            match &self.num {
32                Num::X1 => {
33                    push_directive(tokens, "x1");
34                }
35                Num::X2 => {
36                    push_directive(tokens, "x2");
37                }
38                Num::X4 => {
39                    push_directive(tokens, "x4");
40                }
41            }
42            if self.trans {
43                push_directive(tokens, "trans");
44            }
45            if let Some(ss_0) = self.ss.as_ref() {
46                match ss_0 {
47                    Ss::SharedCta => {
48                        push_directive(tokens, "shared::cta");
49                    }
50                    Ss::Shared => {
51                        push_directive(tokens, "shared");
52                    }
53                }
54            }
55            match &self.type_ {
56                Type::B16 => {
57                    push_directive(tokens, "b16");
58                }
59                Type::B8 => {
60                    push_directive(tokens, "b8");
61                }
62            }
63            self.p.unparse_tokens(tokens);
64            tokens.push(PtxToken::Comma);
65            self.r.unparse_tokens(tokens);
66            tokens.push(PtxToken::Semicolon);
67        }
68    }
69}
ptx_parser/unparser/instruction/stmatrix.rs

ptx_parser/unparser/instruction/
stmatrix.rs
