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            self.unparse_tokens_mode(tokens, false);
21        }
22        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
23            push_opcode(tokens, "stmatrix");
24            push_directive(tokens, "sync");
25            push_directive(tokens, "aligned");
26            match &self.shape {
27                Shape::M16n8 => {
28                    push_directive(tokens, "m16n8");
29                }
30                Shape::M8n8 => {
31                    push_directive(tokens, "m8n8");
32                }
33            }
34            match &self.num {
35                Num::X1 => {
36                    push_directive(tokens, "x1");
37                }
38                Num::X2 => {
39                    push_directive(tokens, "x2");
40                }
41                Num::X4 => {
42                    push_directive(tokens, "x4");
43                }
44            }
45            if self.trans {
46                push_directive(tokens, "trans");
47            }
48            if let Some(ss_0) = self.ss.as_ref() {
49                match ss_0 {
50                    Ss::SharedCta => {
51                        push_directive(tokens, "shared::cta");
52                    }
53                    Ss::Shared => {
54                        push_directive(tokens, "shared");
55                    }
56                }
57            }
58            match &self.type_ {
59                Type::B16 => {
60                    push_directive(tokens, "b16");
61                }
62                Type::B8 => {
63                    push_directive(tokens, "b8");
64                }
65            }
66            if spaced {
67                tokens.push(PtxToken::Space);
68            }
69            self.p.unparse_tokens_mode(tokens, spaced);
70            tokens.push(PtxToken::Comma);
71            if spaced {
72                tokens.push(PtxToken::Space);
73            }
74            self.r.unparse_tokens_mode(tokens, spaced);
75            tokens.push(PtxToken::Semicolon);
76            if spaced {
77                tokens.push(PtxToken::Newline);
78            }
79        }
80    }
81}
ptx_parser/unparser/instruction/stmatrix.rs

ptx_parser/unparser/instruction/
stmatrix.rs
