Add apply_fallback_impl struct

Author: stijnh

include/kernel_float/apply.h

@@ -157,31 +157,53 @@ using default_policy = KERNEL_FLOAT_POLICY;
 
 namespace detail {
 
+//
 template<typename Policy, typename F, size_t N, typename Output, typename... Args>
-struct apply_base_impl {
+struct apply_fallback_impl {
     KERNEL_FLOAT_INLINE static void call(F fun, Output* output, const Args*... args) {
-#pragma unroll
-        for (size_t i = 0; i < N; i++) {
-            output[i] = fun(args[i]...);
-        }
+        static_assert(N > 0, "operation not implemented");
     }
 };
 
+template<typename Policy, typename F, size_t N, typename Output, typename... Args>
+struct apply_base_impl: apply_fallback_impl<Policy, F, N, Output, Args...> {};
+
 template<typename Policy, typename F, size_t N, typename Output, typename... Args>
 struct apply_impl: apply_base_impl<Policy, F, N, Output, Args...> {};
 
+// `fast_policy` falls back to `accurate_policy`
 template<typename F, size_t N, typename Output, typename... Args>
-struct apply_base_impl<fast_policy, F, N, Output, Args...>:
+struct apply_fallback_impl<fast_policy, F, N, Output, Args...>:
     apply_impl<accurate_policy, F, N, Output, Args...> {};
 
+// `approx_policy` falls back to `fast_policy`
 template<typename F, size_t N, typename Output, typename... Args>
-struct apply_base_impl<approx_policy, F, N, Output, Args...>:
+struct apply_fallback_impl<approx_policy, F, N, Output, Args...>:
     apply_impl<fast_policy, F, N, Output, Args...> {};
 
+// `approx_level_policy` falls back to `approx_policy`
 template<int Level, typename F, size_t N, typename Output, typename... Args>
-struct apply_base_impl<approx_level_policy<Level>, F, N, Output, Args...>:
+struct apply_fallback_impl<approx_level_policy<Level>, F, N, Output, Args...>:
     apply_impl<approx_policy, F, N, Output, Args...> {};
 
+template<typename F, typename Output, typename... Args>
+struct invoke_impl {
+    KERNEL_FLOAT_INLINE static Output call(F fun, Args... args) {
+        return fun(args...);
+    }
+};
+
+// Only for `accurate_policy` do we implement `apply_impl`, the others will fall back to `apply_base_impl`.
+template<typename F, size_t N, typename Output, typename... Args>
+struct apply_impl<accurate_policy, F, N, Output, Args...> {
+    KERNEL_FLOAT_INLINE static void call(F fun, Output* output, const Args*... args) {
+#pragma unroll
+        for (size_t i = 0; i < N; i++) {
+            output[i] = invoke_impl<F, Output, Args...>::call(fun, args[i]...);
+        }
+    }
+};
+
 template<typename Policy, typename F, size_t N, typename Output, typename... Args>
 struct map_impl {
     static constexpr size_t packet_size = preferred_vector_size<Output>::value;

include/kernel_float/binops.h

@@ -291,7 +291,7 @@ struct multiply<bool> {
 
 namespace detail {
 template<typename Policy, typename T, size_t N>
-struct apply_impl<Policy, ops::divide<T>, N, T, T, T> {
+struct apply_base_impl<Policy, ops::divide<T>, N, T, T, T> {
     KERNEL_FLOAT_INLINE static void call(ops::divide<T>, T* result, const T* lhs, const T* rhs) {
         T rhs_rcp[N];
 
@@ -301,10 +301,6 @@ struct apply_impl<Policy, ops::divide<T>, N, T, T, T> {
     }
 };
 
-template<typename T, size_t N>
-struct apply_impl<accurate_policy, ops::divide<T>, N, T, T, T>:
-    apply_base_impl<accurate_policy, ops::divide<T>, N, T, T, T> {};
-
 #if KERNEL_FLOAT_IS_DEVICE
 template<>
 struct apply_impl<fast_policy, ops::divide<float>, 1, float, float, float> {
@@ -319,7 +315,7 @@ struct apply_impl<fast_policy, ops::divide<float>, 1, float, float, float> {
 namespace detail {
 // Override `pow` using `log2` and `exp2`
 template<typename Policy, typename T, size_t N>
-struct apply_impl<Policy, ops::pow<T>, N, T, T, T> {
+struct apply_base_impl<Policy, ops::pow<T>, N, T, T, T> {
     KERNEL_FLOAT_INLINE static void call(ops::divide<T>, T* result, const T* lhs, const T* rhs) {
         T lhs_log[N];
         T result_log[N];
@@ -330,10 +326,6 @@ struct apply_impl<Policy, ops::pow<T>, N, T, T, T> {
         apply_impl<Policy, ops::exp2<T>, N, T, T, T>::call({}, result, result_log);
     }
 };
-
-template<typename T, size_t N>
-struct apply_impl<accurate_policy, ops::pow<T>, N, T, T, T>:
-    apply_base_impl<accurate_policy, ops::pow<T>, N, T, T, T> {};
 }  // namespace detail
 
 template<typename L, typename R, typename T = promoted_vector_value_type<L, R>>

include/kernel_float/fp16.h

@@ -22,9 +22,6 @@ namespace kernel_float {
 using half_t = ::__half;
 using half2_t = ::__half2;
 
-using __half = void;
-using __half2 = void;
-
 template<>
 struct preferred_vector_size<half_t> {
     static constexpr size_t value = 2;
@@ -50,7 +47,7 @@ template<>
 struct allow_float_fallback<half_t> {
     static constexpr bool value = true;
 };
-};  // namespace detail
+}  // namespace detail
 
 #if KERNEL_FLOAT_IS_DEVICE
 #define KERNEL_FLOAT_FP16_UNARY_FUN(NAME, FUN1, FUN2)                                              \

include/kernel_float/triops.h

@@ -98,13 +98,13 @@ struct fma {
 }  // namespace ops
 
 namespace detail {
-template<typename Policy, typename T, size_t N>
-struct apply_impl<Policy, ops::fma<T>, N, T, T, T, T> {
+template<typename T, size_t N>
+struct apply_impl<accurate_policy, ops::fma<T>, N, T, T, T, T> {
     KERNEL_FLOAT_INLINE
     static void call(ops::fma<T>, T* output, const T* a, const T* b, const T* c) {
         T temp[N];
-        apply_impl<Policy, ops::multiply<T>, N, T, T, T>::call({}, temp, a, b);
-        apply_impl<Policy, ops::add<T>, N, T, T, T>::call({}, output, temp, c);
+        apply_impl<accurate_policy, ops::multiply<T>, N, T, T, T>::call({}, temp, a, b);
+        apply_impl<accurate_policy, ops::add<T>, N, T, T, T>::call({}, output, temp, c);
     }
 };
 }  // namespace detail