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
70}
71
ptx_parser/unparser/instruction/stmatrix.rs

ptx_parser/unparser/instruction/
stmatrix.rs
