ptx_parser/unparser/instruction/
ldmatrix.rs

1//! Original PTX specification:
2//!
3//! ldmatrix.sync.aligned.shape.num{.trans}{.ss}.type r, [p];
4//! ldmatrix.sync.aligned.m8n16.num{.ss}.dst_fmt.src_fmt        r, [p];
5//! ldmatrix.sync.aligned.m16n16.num.trans{.ss}.dst_fmt.src_fmt r, [p];
6//! .shape   = {.m8n8, .m16n16};
7//! .num     = {.x1, .x2, .x4};
8//! .ss      = {.shared, .shared::cta};
9//! .type    = {.b16, .b8};
10//! .dst_fmt = { .b8x16 };
11//! .src_fmt = { .b6x16_p32, .b4x16_p64 };
12
13#![allow(unused)]
14
15use crate::lexer::PtxToken;
16use crate::unparser::{PtxUnparser, common::*};
17
18pub mod section_0 {
19    use super::*;
20    use crate::r#type::instruction::ldmatrix::section_0::*;
21
22    impl PtxUnparser for LdmatrixSyncAlignedShapeNumTransSsType {
23        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
24            self.unparse_tokens_mode(tokens, false);
25        }
26        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
27            push_opcode(tokens, "ldmatrix");
28            push_directive(tokens, "sync");
29            push_directive(tokens, "aligned");
30            match &self.shape {
31                Shape::M16n16 => {
32                    push_directive(tokens, "m16n16");
33                }
34                Shape::M8n8 => {
35                    push_directive(tokens, "m8n8");
36                }
37            }
38            match &self.num {
39                Num::X1 => {
40                    push_directive(tokens, "x1");
41                }
42                Num::X2 => {
43                    push_directive(tokens, "x2");
44                }
45                Num::X4 => {
46                    push_directive(tokens, "x4");
47                }
48            }
49            if self.trans {
50                push_directive(tokens, "trans");
51            }
52            if let Some(ss_0) = self.ss.as_ref() {
53                match ss_0 {
54                    Ss::SharedCta => {
55                        push_directive(tokens, "shared::cta");
56                    }
57                    Ss::Shared => {
58                        push_directive(tokens, "shared");
59                    }
60                }
61            }
62            match &self.type_ {
63                Type::B16 => {
64                    push_directive(tokens, "b16");
65                }
66                Type::B8 => {
67                    push_directive(tokens, "b8");
68                }
69            }
70            if spaced {
71                tokens.push(PtxToken::Space);
72            }
73            self.r.unparse_tokens_mode(tokens, spaced);
74            tokens.push(PtxToken::Comma);
75            if spaced {
76                tokens.push(PtxToken::Space);
77            }
78            self.p.unparse_tokens_mode(tokens, spaced);
79            tokens.push(PtxToken::Semicolon);
80            if spaced {
81                tokens.push(PtxToken::Newline);
82            }
83        }
84    }
85
86    impl PtxUnparser for LdmatrixSyncAlignedM8n16NumSsDstFmtSrcFmt {
87        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
88            self.unparse_tokens_mode(tokens, false);
89        }
90        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
91            push_opcode(tokens, "ldmatrix");
92            push_directive(tokens, "sync");
93            push_directive(tokens, "aligned");
94            push_directive(tokens, "m8n16");
95            match &self.num {
96                Num::X1 => {
97                    push_directive(tokens, "x1");
98                }
99                Num::X2 => {
100                    push_directive(tokens, "x2");
101                }
102                Num::X4 => {
103                    push_directive(tokens, "x4");
104                }
105            }
106            if let Some(ss_1) = self.ss.as_ref() {
107                match ss_1 {
108                    Ss::SharedCta => {
109                        push_directive(tokens, "shared::cta");
110                    }
111                    Ss::Shared => {
112                        push_directive(tokens, "shared");
113                    }
114                }
115            }
116            match &self.dst_fmt {
117                DstFmt::B8x16 => {
118                    push_directive(tokens, "b8x16");
119                }
120            }
121            match &self.src_fmt {
122                SrcFmt::B6x16P32 => {
123                    push_directive(tokens, "b6x16_p32");
124                }
125                SrcFmt::B4x16P64 => {
126                    push_directive(tokens, "b4x16_p64");
127                }
128            }
129            if spaced {
130                tokens.push(PtxToken::Space);
131            }
132            self.r.unparse_tokens_mode(tokens, spaced);
133            tokens.push(PtxToken::Comma);
134            if spaced {
135                tokens.push(PtxToken::Space);
136            }
137            self.p.unparse_tokens_mode(tokens, spaced);
138            tokens.push(PtxToken::Semicolon);
139            if spaced {
140                tokens.push(PtxToken::Newline);
141            }
142        }
143    }
144
145    impl PtxUnparser for LdmatrixSyncAlignedM16n16NumTransSsDstFmtSrcFmt {
146        fn unparse_tokens(&self, tokens: &mut ::std::vec::Vec<PtxToken>) {
147            self.unparse_tokens_mode(tokens, false);
148        }
149        fn unparse_tokens_mode(&self, tokens: &mut ::std::vec::Vec<PtxToken>, spaced: bool) {
150            push_opcode(tokens, "ldmatrix");
151            push_directive(tokens, "sync");
152            push_directive(tokens, "aligned");
153            push_directive(tokens, "m16n16");
154            match &self.num {
155                Num::X1 => {
156                    push_directive(tokens, "x1");
157                }
158                Num::X2 => {
159                    push_directive(tokens, "x2");
160                }
161                Num::X4 => {
162                    push_directive(tokens, "x4");
163                }
164            }
165            push_directive(tokens, "trans");
166            if let Some(ss_2) = self.ss.as_ref() {
167                match ss_2 {
168                    Ss::SharedCta => {
169                        push_directive(tokens, "shared::cta");
170                    }
171                    Ss::Shared => {
172                        push_directive(tokens, "shared");
173                    }
174                }
175            }
176            match &self.dst_fmt {
177                DstFmt::B8x16 => {
178                    push_directive(tokens, "b8x16");
179                }
180            }
181            match &self.src_fmt {
182                SrcFmt::B6x16P32 => {
183                    push_directive(tokens, "b6x16_p32");
184                }
185                SrcFmt::B4x16P64 => {
186                    push_directive(tokens, "b4x16_p64");
187                }
188            }
189            if spaced {
190                tokens.push(PtxToken::Space);
191            }
192            self.r.unparse_tokens_mode(tokens, spaced);
193            tokens.push(PtxToken::Comma);
194            if spaced {
195                tokens.push(PtxToken::Space);
196            }
197            self.p.unparse_tokens_mode(tokens, spaced);
198            tokens.push(PtxToken::Semicolon);
199            if spaced {
200                tokens.push(PtxToken::Newline);
201            }
202        }
203    }
204}
ptx_parser/unparser/instruction/ldmatrix.rs

ptx_parser/unparser/instruction/
ldmatrix.rs
