Skip to main content

cubecl_cpp/shared/
unary.rs

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            // Dialects without a half prefix (e.g. Metal) reuse the float-taking
109            // function and lack `bfloat` intrinsics, so bfloat must round-trip through
110            // float. Dialects with a prefix (e.g. CUDA's `h`) have native intrinsics.
111            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                // Small ints cause issues around auto-promotion on AMD, so use explicit cast on out
124                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        // Bitwise negation may widen, so use explicit cast
372        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        // Cast only when necessary.
420        if elem != input.elem() {
421            match elem {
422                Elem::TF32 => write!(f, "nvcuda::wmma::__float_to_tf32({input})"),
423                // A direct construction between the two half types is
424                // ambiguous on HIP — `__hip_bfloat16` exposes a dozen implicit
425                // conversion operators, so `__half(bf16)` (and the reverse)
426                // fails to compile. Route half<->half through `float`, which
427                // is a single unambiguous conversion on every dialect.
428                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
443/// The two scalar half-precision float types, whose direct
444/// interconstruction is ambiguous on HIP.
445fn 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    // Math functions prefix (no leading underscores)
451    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
465// `isnan` / `isinf` are defined for cuda/hip/metal with same prefixes for half/bf16 on cuda/hip
466pub 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        // Format unary function name based on *input* elem dtype
475        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        // Format unary function name based on *input* elem dtype
493        let elem = input.elem();
494        write!(f, "{}({input})", elem_function_name("isinf", elem))
495    }
496
497    fn can_optimize() -> bool {
498        true
499    }
500}