1use super::{Component, Dialect, Elem, FmtLeft, Variable};
2use std::fmt::Display;
3
4pub trait Unary<D: Dialect> {
5 fn format(
6 f: &mut std::fmt::Formatter<'_>,
7 input: &Variable<D>,
8 out: &Variable<D>,
9 ) -> std::fmt::Result {
10 let out_item = out.item();
11
12 if out_item.vectorization == 1 {
13 write!(f, "{} = ", out.fmt_left())?;
14 Self::format_scalar(f, *input, out_item.elem)?;
15 f.write_str(";\n")
16 } else {
17 Self::unroll_vec(f, input, out, out_item.elem, out_item.vectorization)
18 }
19 }
20
21 fn format_scalar<Input: Component<D>>(
22 f: &mut std::fmt::Formatter<'_>,
23 input: Input,
24 out_elem: Elem<D>,
25 ) -> std::fmt::Result;
26
27 fn unroll_vec(
28 f: &mut std::fmt::Formatter<'_>,
29 input: &Variable<D>,
30 out: &Variable<D>,
31 out_elem: Elem<D>,
32 index: usize,
33 ) -> std::fmt::Result {
34 let mut write_op = |index, out_elem, input: &Variable<D>, out: &Variable<D>| {
35 let out_item = out.item();
36 let out = out.fmt_left();
37 writeln!(f, "{out} = {out_item}{{")?;
38
39 for i in 0..index {
40 let inputi = input.index(i);
41
42 Self::format_scalar(f, inputi, out_elem)?;
43 f.write_str(",")?;
44 }
45
46 f.write_str("};\n")
47 };
48
49 if Self::can_optimize() {
50 let optimized = Variable::optimized_args([*input, *out]);
51 let [input, out_optimized] = optimized.args;
52
53 let item_out_original = out.item();
54 let item_out_optimized = out_optimized.item();
55
56 let (index, out_elem) = match optimized.optimization_factor {
57 Some(factor) => (index / factor, out_optimized.elem()),
58 None => (index, out_elem),
59 };
60
61 if item_out_original != item_out_optimized {
62 let out_tmp = Variable::tmp(item_out_optimized);
63
64 write_op(index, out_elem, &input, &out_tmp)?;
65 let qualifier = out.const_qualifier();
66 let addr_space = D::address_space_for_variable(out);
67 let out_fmt = out.fmt_left();
68 writeln!(
69 f,
70 "{out_fmt} = reinterpret_cast<{addr_space}{item_out_original}{qualifier}&>({out_tmp});\n"
71 )
72 } else {
73 write_op(index, out_elem, &input, &out_optimized)
74 }
75 } else {
76 write_op(index, out_elem, input, out)
77 }
78 }
79
80 fn can_optimize() -> bool {
81 true
82 }
83}
84
85pub trait FunctionFmt<D: Dialect> {
86 fn base_function_name() -> &'static str;
87 fn function_name(elem: Elem<D>) -> String {
88 if Self::half_support() {
89 let prefix = match elem {
90 Elem::F16 | Elem::BF16 => D::compile_instruction_half_function_name_prefix(),
91 Elem::F16x2 | Elem::BF16x2 => D::compile_instruction_half2_function_name_prefix(),
92 _ => "",
93 };
94 format!("{prefix}{}", Self::base_function_name())
95 } else {
96 Self::base_function_name().into()
97 }
98 }
99 fn format_unary<Input: Display>(
100 f: &mut std::fmt::Formatter<'_>,
101 input: Input,
102 elem: Elem<D>,
103 ) -> std::fmt::Result {
104 if Self::half_support() {
105 write!(f, "{}({input})", Self::function_name(elem))
106 } else {
107 match elem {
108 Elem::F16 | Elem::F16x2 | Elem::BF16 | Elem::BF16x2 => {
109 write!(f, "{}({}(float({input})))", elem, Self::function_name(elem))
110 }
111 _ => write!(f, "{}({input})", Self::function_name(elem)),
112 }
113 }
114 }
115
116 fn half_support() -> bool;
117}
118
119macro_rules! function {
120 ($name:ident, $func:expr) => {
121 function!($name, $func, true);
122 };
123 ($name:ident, $func:expr, $half_support:expr) => {
124 pub struct $name;
125
126 impl<D: Dialect> FunctionFmt<D> for $name {
127 fn base_function_name() -> &'static str {
128 $func
129 }
130 fn half_support() -> bool {
131 $half_support
132 }
133 }
134
135 impl<D: Dialect> Unary<D> for $name {
136 fn format_scalar<Input: Display>(
137 f: &mut std::fmt::Formatter<'_>,
138 input: Input,
139 elem: Elem<D>,
140 ) -> std::fmt::Result {
141 Self::format_unary(f, input, elem)
142 }
143
144 fn can_optimize() -> bool {
145 $half_support
146 }
147 }
148 };
149}
150
151function!(Log, "log");
152function!(FastLog, "__logf", false);
153function!(Cos, "cos");
154function!(FastCos, "__cosf", false);
155function!(Sin, "sin");
156function!(FastSin, "__sinf", false);
157function!(Sqrt, "sqrt");
158function!(InverseSqrt, "rsqrt");
159function!(FastSqrt, "__fsqrt_rn", false);
160function!(FastInverseSqrt, "__frsqrt_rn", false);
161function!(Exp, "exp");
162function!(FastExp, "__expf", false);
163function!(Ceil, "ceil");
164function!(Trunc, "trunc");
165function!(Floor, "floor");
166function!(Round, "rint");
167function!(FastRecip, "__frcp_rn", false);
168function!(FastTanh, "__tanhf", false);
169
170function!(Erf, "erf", false);
171function!(Abs, "abs", false);
172
173pub struct Log1p;
174
175impl<D: Dialect> Unary<D> for Log1p {
176 fn format_scalar<Input: Component<D>>(
177 f: &mut std::fmt::Formatter<'_>,
178 input: Input,
179 _out_elem: Elem<D>,
180 ) -> std::fmt::Result {
181 D::compile_instruction_log1p_scalar(f, input)
182 }
183
184 fn can_optimize() -> bool {
185 false
186 }
187}
188
189pub struct Tanh;
190
191impl<D: Dialect> Unary<D> for Tanh {
192 fn format_scalar<Input: Component<D>>(
193 f: &mut std::fmt::Formatter<'_>,
194 input: Input,
195 _out_elem: Elem<D>,
196 ) -> std::fmt::Result {
197 D::compile_instruction_tanh_scalar(f, input)
198 }
199
200 fn can_optimize() -> bool {
201 false
202 }
203}
204
205pub fn zero_extend<D: Dialect>(input: impl Component<D>) -> String {
206 match input.elem() {
207 Elem::I8 => format!("{}({}({input}))", Elem::<D>::U32, Elem::<D>::U8),
208 Elem::I16 => format!("{}({}({input}))", Elem::<D>::U32, Elem::<D>::U16),
209 Elem::U8 => format!("{}({input})", Elem::<D>::U32),
210 Elem::U16 => format!("{}({input})", Elem::<D>::U32),
211 _ => unreachable!("zero extend only supports integer < 32 bits"),
212 }
213}
214
215pub struct CountBits;
216
217impl<D: Dialect> Unary<D> for CountBits {
218 fn format_scalar<Input: Component<D>>(
219 f: &mut std::fmt::Formatter<'_>,
220 input: Input,
221 elem: Elem<D>,
222 ) -> std::fmt::Result {
223 D::compile_instruction_popcount_scalar(f, input, elem)
224 }
225}
226
227pub struct ReverseBits;
228
229impl<D: Dialect> Unary<D> for ReverseBits {
230 fn format_scalar<Input: Component<D>>(
231 f: &mut std::fmt::Formatter<'_>,
232 input: Input,
233 elem: Elem<D>,
234 ) -> std::fmt::Result {
235 D::compile_instruction_reverse_bits_scalar(f, input, elem)
236 }
237}
238
239pub struct LeadingZeros;
240
241impl<D: Dialect> Unary<D> for LeadingZeros {
242 fn format_scalar<Input: Component<D>>(
243 f: &mut std::fmt::Formatter<'_>,
244 input: Input,
245 elem: Elem<D>,
246 ) -> std::fmt::Result {
247 D::compile_instruction_leading_zeros_scalar(f, input, elem)
248 }
249}
250
251pub struct FindFirstSet;
252
253impl<D: Dialect> Unary<D> for FindFirstSet {
254 fn format_scalar<Input: Component<D>>(
255 f: &mut std::fmt::Formatter<'_>,
256 input: Input,
257 out_elem: Elem<D>,
258 ) -> std::fmt::Result {
259 D::compile_instruction_find_first_set(f, input, out_elem)
260 }
261}
262
263pub struct BitwiseNot;
264
265impl<D: Dialect> Unary<D> for BitwiseNot {
266 fn format_scalar<Input>(
267 f: &mut std::fmt::Formatter<'_>,
268 input: Input,
269 _out_elem: Elem<D>,
270 ) -> std::fmt::Result
271 where
272 Input: Component<D>,
273 {
274 write!(f, "~{input}")
275 }
276}
277
278pub struct Not;
279
280impl<D: Dialect> Unary<D> for Not {
281 fn format_scalar<Input>(
282 f: &mut std::fmt::Formatter<'_>,
283 input: Input,
284 _out_elem: Elem<D>,
285 ) -> std::fmt::Result
286 where
287 Input: Component<D>,
288 {
289 write!(f, "!{input}")
290 }
291}
292
293pub struct Assign;
294
295impl<D: Dialect> Unary<D> for Assign {
296 fn format(
297 f: &mut std::fmt::Formatter<'_>,
298 input: &Variable<D>,
299 out: &Variable<D>,
300 ) -> std::fmt::Result {
301 let item = out.item();
302
303 if item.vectorization == 1 || input.item() == item {
304 write!(f, "{} = ", out.fmt_left())?;
305 Self::format_scalar(f, *input, item.elem)?;
306 f.write_str(";\n")
307 } else {
308 Self::unroll_vec(f, input, out, item.elem, item.vectorization)
309 }
310 }
311
312 fn format_scalar<Input>(
313 f: &mut std::fmt::Formatter<'_>,
314 input: Input,
315 elem: Elem<D>,
316 ) -> std::fmt::Result
317 where
318 Input: Component<D>,
319 {
320 if elem != input.elem() {
322 match elem {
323 Elem::TF32 => write!(f, "nvcuda::wmma::__float_to_tf32({input})"),
324 elem => write!(f, "{elem}({input})"),
325 }
326 } else {
327 write!(f, "{input}")
328 }
329 }
330}
331
332fn elem_function_name<D: Dialect>(base_name: &'static str, elem: Elem<D>) -> String {
333 let prefix = match elem {
335 Elem::F16 | Elem::BF16 => D::compile_instruction_half_function_name_prefix(),
336 Elem::F16x2 | Elem::BF16x2 => D::compile_instruction_half2_function_name_prefix(),
337 _ => "",
338 };
339 if prefix.is_empty() {
340 base_name.to_string()
341 } else if prefix == "h" || prefix == "h2" {
342 format!("__{prefix}{base_name}")
343 } else {
344 panic!("Unknown prefix '{prefix}'");
345 }
346}
347
348pub struct IsNan;
350
351impl<D: Dialect> Unary<D> for IsNan {
352 fn format_scalar<Input: Component<D>>(
353 f: &mut std::fmt::Formatter<'_>,
354 input: Input,
355 _elem: Elem<D>,
356 ) -> std::fmt::Result {
357 let elem = input.elem();
359 write!(f, "{}({input})", elem_function_name("isnan", elem))
360 }
361
362 fn can_optimize() -> bool {
363 true
364 }
365}
366
367pub struct IsInf;
368
369impl<D: Dialect> Unary<D> for IsInf {
370 fn format_scalar<Input: Component<D>>(
371 f: &mut std::fmt::Formatter<'_>,
372 input: Input,
373 _elem: Elem<D>,
374 ) -> std::fmt::Result {
375 let elem = input.elem();
377 write!(f, "{}({input})", elem_function_name("isinf", elem))
378 }
379
380 fn can_optimize() -> bool {
381 true
382 }
383}