1use super::{Component, Dialect, Elem, FmtLeft, Value};
2use std::fmt::Display;
3
4pub trait Unary<D: Dialect> {
5 fn format(
6 f: &mut std::fmt::Formatter<'_>,
7 input: &Value<D>,
8 out: &Value<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: &Value<D>,
30 out: &Value<D>,
31 out_elem: Elem<D>,
32 index: usize,
33 ) -> std::fmt::Result {
34 let mut write_op = |index, out_elem, input: &Value<D>, out: &Value<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 = Value::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 = Value::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_value(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!(
95 "{prefix}{}",
96 D::compile_fast_math_function_name(Self::base_function_name())
97 )
98 } else {
99 D::compile_fast_math_function_name(Self::base_function_name()).into()
100 }
101 }
102 fn format_unary<Input: Display>(
103 f: &mut std::fmt::Formatter<'_>,
104 input: Input,
105 elem: Elem<D>,
106 ) -> std::fmt::Result {
107 if Self::half_support() {
108 let no_half_prefix = D::compile_instruction_half_function_name_prefix().is_empty();
112 match elem {
113 Elem::BF16 | Elem::BF16x2 if no_half_prefix => {
114 write!(f, "{}({}(float({input})))", elem, Self::function_name(elem))
115 }
116 _ => write!(f, "{}({input})", Self::function_name(elem)),
117 }
118 } else {
119 match elem {
120 Elem::F16 | Elem::F16x2 | Elem::BF16 | Elem::BF16x2 => {
121 write!(f, "{}({}(float({input})))", elem, Self::function_name(elem))
122 }
123 Elem::U16 | Elem::U8 | Elem::I16 | Elem::I8 => {
125 write!(f, "{elem}({}({input}))", Self::function_name(elem))
126 }
127 _ => write!(f, "{}({input})", Self::function_name(elem)),
128 }
129 }
130 }
131
132 fn half_support() -> bool;
133}
134
135macro_rules! function {
136 ($name:ident, $func:expr) => {
137 function!($name, $func, true);
138 };
139 ($name:ident, $func:expr, $half_support:expr) => {
140 pub struct $name;
141
142 impl<D: Dialect> FunctionFmt<D> for $name {
143 fn base_function_name() -> &'static str {
144 $func
145 }
146 fn half_support() -> bool {
147 $half_support
148 }
149 }
150
151 impl<D: Dialect> Unary<D> for $name {
152 fn format_scalar<Input: Display>(
153 f: &mut std::fmt::Formatter<'_>,
154 input: Input,
155 elem: Elem<D>,
156 ) -> std::fmt::Result {
157 Self::format_unary(f, input, elem)
158 }
159
160 fn can_optimize() -> bool {
161 $half_support
162 }
163 }
164 };
165}
166
167function!(Log, "log");
168function!(FastLog, "__logf", false);
169function!(Sin, "sin");
170function!(Cos, "cos");
171function!(Tan, "tan", false);
172function!(Sinh, "sinh", false);
173function!(Cosh, "cosh", false);
174function!(ArcCos, "acos", false);
175function!(ArcSin, "asin", false);
176function!(ArcTan, "atan", false);
177function!(ArcSinh, "asinh", false);
178function!(ArcCosh, "acosh", false);
179function!(ArcTanh, "atanh", false);
180function!(FastSin, "__sinf", false);
181function!(FastCos, "__cosf", false);
182function!(Sqrt, "sqrt");
183function!(InverseSqrt, "rsqrt");
184function!(FastSqrt, "__fsqrt_rn", false);
185function!(FastInverseSqrt, "__frsqrt_rn", false);
186function!(Exp, "exp");
187function!(FastExp, "__expf", false);
188function!(Ceil, "ceil");
189function!(Trunc, "trunc");
190function!(Floor, "floor");
191function!(Round, "rint");
192function!(FastRecip, "__frcp_rn", false);
193function!(FastTanh, "__tanhf", false);
194
195function!(Erf, "erf", false);
196function!(Abs, "abs", false);
197
198pub struct Neg;
199
200impl<D: Dialect> Unary<D> for Neg {
201 fn format_scalar<Input: Component<D>>(
202 f: &mut std::fmt::Formatter<'_>,
203 input: Input,
204 _out_elem: Elem<D>,
205 ) -> std::fmt::Result {
206 writeln!(f, "-{}", input)
207 }
208}
209
210pub struct Log1p;
211
212impl<D: Dialect> Unary<D> for Log1p {
213 fn format_scalar<Input: Component<D>>(
214 f: &mut std::fmt::Formatter<'_>,
215 input: Input,
216 _out_elem: Elem<D>,
217 ) -> std::fmt::Result {
218 D::compile_instruction_log1p_scalar(f, input)
219 }
220
221 fn can_optimize() -> bool {
222 false
223 }
224}
225
226pub struct Expm1;
227
228impl<D: Dialect> Unary<D> for Expm1 {
229 fn format_scalar<Input: Component<D>>(
230 f: &mut std::fmt::Formatter<'_>,
231 input: Input,
232 _out_elem: Elem<D>,
233 ) -> std::fmt::Result {
234 D::compile_instruction_expm1_scalar(f, input)
235 }
236
237 fn can_optimize() -> bool {
238 false
239 }
240}
241
242pub struct Tanh;
243
244impl<D: Dialect> Unary<D> for Tanh {
245 fn format_scalar<Input: Component<D>>(
246 f: &mut std::fmt::Formatter<'_>,
247 input: Input,
248 _out_elem: Elem<D>,
249 ) -> std::fmt::Result {
250 D::compile_instruction_tanh_scalar(f, input)
251 }
252
253 fn can_optimize() -> bool {
254 false
255 }
256}
257
258pub struct Degrees;
259
260impl<D: Dialect> Unary<D> for Degrees {
261 fn format_scalar<Input: Component<D>>(
262 f: &mut std::fmt::Formatter<'_>,
263 input: Input,
264 elem: Elem<D>,
265 ) -> std::fmt::Result {
266 write!(f, "{input}*{elem}(57.29577951308232f)")
267 }
268
269 fn can_optimize() -> bool {
270 false
271 }
272}
273
274pub struct Radians;
275
276impl<D: Dialect> Unary<D> for Radians {
277 fn format_scalar<Input: Component<D>>(
278 f: &mut std::fmt::Formatter<'_>,
279 input: Input,
280 elem: Elem<D>,
281 ) -> std::fmt::Result {
282 write!(f, "{input}*{elem}(0.017453292519943295f)")
283 }
284
285 fn can_optimize() -> bool {
286 false
287 }
288}
289
290pub fn zero_extend<D: Dialect>(input: impl Component<D>) -> String {
291 match input.elem() {
292 Elem::I8 => format!("{}({}({input}))", Elem::<D>::U32, Elem::<D>::U8),
293 Elem::I16 => format!("{}({}({input}))", Elem::<D>::U32, Elem::<D>::U16),
294 Elem::U8 => format!("{}({input})", Elem::<D>::U32),
295 Elem::U16 => format!("{}({input})", Elem::<D>::U32),
296 _ => unreachable!("zero extend only supports integer < 32 bits"),
297 }
298}
299
300pub struct CountBits;
301
302impl<D: Dialect> Unary<D> for CountBits {
303 fn format_scalar<Input: Component<D>>(
304 f: &mut std::fmt::Formatter<'_>,
305 input: Input,
306 elem: Elem<D>,
307 ) -> std::fmt::Result {
308 D::compile_instruction_popcount_scalar(f, input, elem)
309 }
310}
311
312pub struct ReverseBits;
313
314impl<D: Dialect> Unary<D> for ReverseBits {
315 fn format_scalar<Input: Component<D>>(
316 f: &mut std::fmt::Formatter<'_>,
317 input: Input,
318 elem: Elem<D>,
319 ) -> std::fmt::Result {
320 D::compile_instruction_reverse_bits_scalar(f, input, elem)
321 }
322}
323
324pub struct LeadingZeros;
325
326impl<D: Dialect> Unary<D> for LeadingZeros {
327 fn format_scalar<Input: Component<D>>(
328 f: &mut std::fmt::Formatter<'_>,
329 input: Input,
330 elem: Elem<D>,
331 ) -> std::fmt::Result {
332 D::compile_instruction_leading_zeros_scalar(f, input, elem)
333 }
334}
335
336pub struct TrailingZeros;
337
338impl<D: Dialect> Unary<D> for TrailingZeros {
339 fn format_scalar<Input: Component<D>>(
340 f: &mut std::fmt::Formatter<'_>,
341 input: Input,
342 elem: Elem<D>,
343 ) -> std::fmt::Result {
344 D::compile_instruction_trailing_zeros_scalar(f, input, elem)
345 }
346}
347
348pub struct FindFirstSet;
349
350impl<D: Dialect> Unary<D> for FindFirstSet {
351 fn format_scalar<Input: Component<D>>(
352 f: &mut std::fmt::Formatter<'_>,
353 input: Input,
354 out_elem: Elem<D>,
355 ) -> std::fmt::Result {
356 D::compile_instruction_find_first_set(f, input, out_elem)
357 }
358}
359
360pub struct BitwiseNot;
361
362impl<D: Dialect> Unary<D> for BitwiseNot {
363 fn format_scalar<Input>(
364 f: &mut std::fmt::Formatter<'_>,
365 input: Input,
366 out_elem: Elem<D>,
367 ) -> std::fmt::Result
368 where
369 Input: Component<D>,
370 {
371 write!(f, "{out_elem}(~{input})")
373 }
374}
375
376pub struct Not;
377
378impl<D: Dialect> Unary<D> for Not {
379 fn format_scalar<Input>(
380 f: &mut std::fmt::Formatter<'_>,
381 input: Input,
382 _out_elem: Elem<D>,
383 ) -> std::fmt::Result
384 where
385 Input: Component<D>,
386 {
387 write!(f, "!{input}")
388 }
389}
390
391pub struct Assign;
392
393impl<D: Dialect> Unary<D> for Assign {
394 fn format(
395 f: &mut std::fmt::Formatter<'_>,
396 input: &Value<D>,
397 out: &Value<D>,
398 ) -> std::fmt::Result {
399 let item = out.item();
400 let item = item.value_ty();
401
402 if item.vectorization() == 1 || input.item().value_ty() == item {
403 write!(f, "{} = ", out.fmt_left())?;
404 Self::format_scalar(f, *input, *item.elem())?;
405 f.write_str(";\n")
406 } else {
407 Self::unroll_vec(f, input, out, *item.elem(), item.vectorization())
408 }
409 }
410
411 fn format_scalar<Input>(
412 f: &mut std::fmt::Formatter<'_>,
413 input: Input,
414 elem: Elem<D>,
415 ) -> std::fmt::Result
416 where
417 Input: Component<D>,
418 {
419 if elem != input.elem() {
421 match elem {
422 Elem::TF32 => write!(f, "nvcuda::wmma::__float_to_tf32({input})"),
423 elem if is_half(elem) && is_half(input.elem()) => {
429 write!(f, "{elem}(float({input}))")
430 }
431 elem => write!(f, "{elem}({input})"),
432 }
433 } else {
434 write!(f, "{input}")
435 }
436 }
437
438 fn can_optimize() -> bool {
439 false
440 }
441}
442
443fn is_half<D: Dialect>(elem: Elem<D>) -> bool {
446 matches!(elem, Elem::F16 | Elem::BF16)
447}
448
449fn elem_function_name<D: Dialect>(base_name: &'static str, elem: Elem<D>) -> String {
450 let prefix = match elem {
452 Elem::F16 | Elem::BF16 => D::compile_instruction_half_function_name_prefix(),
453 Elem::F16x2 | Elem::BF16x2 => D::compile_instruction_half2_function_name_prefix(),
454 _ => "",
455 };
456 if prefix.is_empty() {
457 base_name.to_string()
458 } else if prefix == "h" || prefix == "h2" {
459 format!("__{prefix}{base_name}")
460 } else {
461 panic!("Unknown prefix '{prefix}'");
462 }
463}
464
465pub struct IsNan;
467
468impl<D: Dialect> Unary<D> for IsNan {
469 fn format_scalar<Input: Component<D>>(
470 f: &mut std::fmt::Formatter<'_>,
471 input: Input,
472 _elem: Elem<D>,
473 ) -> std::fmt::Result {
474 let elem = input.elem();
476 write!(f, "{}({input})", elem_function_name("isnan", elem))
477 }
478
479 fn can_optimize() -> bool {
480 true
481 }
482}
483
484pub struct IsInf;
485
486impl<D: Dialect> Unary<D> for IsInf {
487 fn format_scalar<Input: Component<D>>(
488 f: &mut std::fmt::Formatter<'_>,
489 input: Input,
490 _elem: Elem<D>,
491 ) -> std::fmt::Result {
492 let elem = input.elem();
494 write!(f, "{}({input})", elem_function_name("isinf", elem))
495 }
496
497 fn can_optimize() -> bool {
498 true
499 }
500}