Fix incorrect definition of FMA

Author: stijnh

include/kernel_float/triops.h

@@ -92,7 +92,7 @@ namespace ops {
 template<typename T>
 struct fma {
     KERNEL_FLOAT_INLINE T operator()(T a, T b, T c) {
-        return a + b * c;
+        return a * b + c;
     }
 };
 

tests/constant.cu

@@ -0,0 +1,30 @@
+#include "common.h"
+
+struct triops_tests {
+    template<typename T, size_t... I, size_t N = sizeof...(I)>
+    __host__ __device__ void operator()(generator<T> gen, std::index_sequence<I...>) {
+        T x[N] = {gen.next(I)...};
+        T y[N] = {gen.next(I)...};
+        T z[N] = {gen.next(I)...};
+
+        kf::vec<T, N> a = {x[I]...};
+        kf::vec<T, N> b = {y[I]...};
+        kf::vec<T, N> c = {z[I]...};
+
+        kf::vec<T, N> answer = kf::where(a, b, c);
+        ASSERT_EQ_ALL(answer[I], bool(x[I]) ? y[I] : z[I]);
+
+        answer = kf::where(a, b);
+        ASSERT_EQ_ALL(answer[I], bool(x[I]) ? y[I] : T());
+
+        answer = kf::where(a);
+        ASSERT_EQ_ALL(answer[I], T(bool(x[I])));
+
+        answer = kf::fma(a, b, c);
+        ASSERT_EQ_ALL(answer[I], x[I] * y[I] + z[I]);
+
+    }
+};
+
+REGISTER_TEST_CASE("ternary operators", triops_tests, int, float, double)
+REGISTER_TEST_CASE_GPU("ternary operators", triops_tests, __half, __nv_bfloat16)

tests/triops.cu

@@ -0,0 +1,137 @@
+#include "common.h"
+
+struct binops_tests {
+    template<typename T, size_t... I, size_t N = sizeof...(I)>
+    __host__ __device__ void operator()(generator<T> gen, std::index_sequence<I...>) {
+        T x[N] = {gen.next(I)...};
+        T y[N] = {gen.next(I)...};
+
+        kf::vec<T, N> a = {x[I]...};
+        kf::vec<T, N> b = {y[I]...};
+        kf::vec<T, N> c;
+
+        // Arithmetic
+        c = a + b;
+        ASSERT(equals(T(x[I] + y[I]), c[I]) && ...);
+
+        c = a - b;
+        ASSERT(equals(T(x[I] - y[I]), c[I]) && ...);
+
+        c = a * b;
+        ASSERT(equals(T(x[I] * y[I]), c[I]) && ...);
+
+        // Results in division by zero
+        //        c = a / b;
+        //        ASSERT(equals(T(x[I] / y[I]), c[I]) && ...);
+
+        // Results in division by zero
+        //        c = a % b;
+        //        ASSERT(equals(T(x[I] % y[I]), c[I]) && ...);
+
+        // Comparison
+        c = a < b;
+        ASSERT(equals(T(x[I] < y[I]), c[I]) && ...);
+
+        c = a > b;
+        ASSERT(equals(T(x[I] > y[I]), c[I]) && ...);
+
+        c = a <= b;
+        ASSERT(equals(T(x[I] <= y[I]), c[I]) && ...);
+
+        c = a >= b;
+        ASSERT(equals(T(x[I] >= y[I]), c[I]) && ...);
+
+        c = a == b;
+        ASSERT(equals(T(x[I] == y[I]), c[I]) && ...);
+
+        c = a != b;
+        ASSERT(equals(T(x[I] != y[I]), c[I]) && ...);
+
+        // Assignment
+        c = a;
+        c += b;
+        ASSERT(equals(T(x[I] + y[I]), c[I]) && ...);
+
+        c = a;
+        c -= b;
+        ASSERT(equals(T(x[I] - y[I]), c[I]) && ...);
+
+        c = a;
+        c *= b;
+        ASSERT(equals(T(x[I] * y[I]), c[I]) && ...);
+    }
+};
+
+REGISTER_TEST_CASE("binary operators", binops_tests, bool, int, float, double)
+REGISTER_TEST_CASE_GPU("binary operators", binops_tests, __half, __nv_bfloat16)
+
+struct binops_float_tests {
+    template<typename T, size_t... I, size_t N = sizeof...(I)>
+    __host__ __device__ void operator()(generator<T> gen, std::index_sequence<I...>) {
+        T x[N] = {gen.next(I)...};
+        T y[N] = {gen.next(I)...};
+
+        kf::vec<T, N> a = {x[I]...};
+        kf::vec<T, N> b = {y[I]...};
+        kf::vec<T, N> c;
+
+        c = a / b;
+        ASSERT(equals(T(x[I] / y[I]), c[I]) && ...);
+
+        // remainder is not support for fp16
+        if constexpr (is_none_of<T, __half, __nv_bfloat16>) {
+            //            c = a % b;
+            //            ASSERT(equals(T(fmod(x[I], y[I])), c[I]) && ...);
+        }
+    }
+};
+
+REGISTER_TEST_CASE("binary float operators", binops_float_tests, float, double)
+REGISTER_TEST_CASE_GPU("binary float operators", binops_float_tests, __half, __nv_bfloat16)
+
+struct minmax_tests {
+    template<typename T, size_t... I, size_t N = sizeof...(I)>
+    __host__ __device__ void operator()(generator<T> gen, std::index_sequence<I...>) {
+        T x[N] = {gen.next(I)...};
+        T y[N] = {gen.next(I)...};
+
+        kf::vec<T, N> a = {x[I]...};
+        kf::vec<T, N> b = {y[I]...};
+
+        kf::vec<T, N> lo = min(a, b);
+        kf::vec<T, N> hi = max(a, b);
+
+        if constexpr (is_one_of<T, double>) {
+            ASSERT(equals(fmin(a[I], b[I]), lo[I]) && ...);
+            ASSERT(equals(fmax(a[I], b[I]), hi[I]) && ...);
+        } else if constexpr (is_one_of<T, float>) {
+            ASSERT(equals(fminf(a[I], b[I]), lo[I]) && ...);
+            ASSERT(equals(fmaxf(a[I], b[I]), hi[I]) && ...);
+        } else if constexpr (is_one_of<T, __half, __nv_bfloat16>) {
+            ASSERT(equals(__hmin(a[I], b[I]), lo[I]) && ...);
+            ASSERT(equals(__hmax(a[I], b[I]), hi[I]) && ...);
+        } else {
+            ASSERT(equals(x[I] < y[I] ? x[I] : y[I], lo[I]) && ...);
+            ASSERT(equals(x[I] < y[I] ? y[I] : x[I], hi[I]) && ...);
+        }
+    }
+};
+
+REGISTER_TEST_CASE("min/max functions", minmax_tests, bool, int, float, double)
+REGISTER_TEST_CASE_GPU("min/max functions", minmax_tests, __half, __nv_bfloat16)
+
+struct cross_test {
+    template<typename T>
+    __host__ __device__ void operator()(generator<T> gen) {
+        kf::vec<T, 3> a = {1, 2, 3};
+        kf::vec<T, 3> b = {4, 5, 6};
+        kf::vec<T, 3> c = cross(a, b);
+
+        ASSERT(c[0] == T(-3));
+        ASSERT(c[1] == T(6));
+        ASSERT(c[2] == T(-3));
+    }
+};
+
+REGISTER_TEST_CASE("cross product", cross_test, float, double)
+REGISTER_TEST_CASE_GPU("cross product", cross_test, __half, __nv_bfloat16)