Format

caa488a9 · Paul · 33716162 · caa488a9
Commit caa488a9 authored Jul 06, 2022 by Paul
Hide whitespace changes
Inline Side-by-side

Showing with 25 additions and 10 deletions

src/targets/gpu/kernels/include/migraphx/kernels/reduce.hpp src/targets/gpu/kernels/include/migraphx/kernels/reduce.hpp +25 -10

No files found.
--- a/src/targets/gpu/kernels/include/migraphx/kernels/reduce.hpp
+++ b/src/targets/gpu/kernels/include/migraphx/kernels/reduce.hpp
@@ -79,16 +79,31 @@ __device__ void dpp_reduce(T& in, Op op)
 #endif
 // NOLINTNEXTLINE
-#define MIGRAPHX_DPP_REDUCE(op, prefix)                                                            \
+#define MIGRAPHX_DPP_REDUCE(op, prefix)                    \
-    __device__ inline void dpp_reduce(double& x, op) { MIGRAPHX_DPP_REDUCE_ASM(x, v_##prefix##_f64); } \
+    __device__ inline void dpp_reduce(double& x, op)       \
-    __device__ inline void dpp_reduce(float& x, op) { MIGRAPHX_DPP_REDUCE_ASM(x, v_##prefix##_f32); }  \
+    {                                                      \
-    __device__ inline void dpp_reduce(half& x, op) { MIGRAPHX_DPP_REDUCE_ASM(x, v_##prefix##_f16); }   \
+        MIGRAPHX_DPP_REDUCE_ASM(x, v_##prefix##_f64);      \
-    __device__ inline void dpp_reduce(int32_t& x, op)                                              \
+    }                                                      \
-    {                                                                                              \
+    __device__ inline void dpp_reduce(float& x, op)        \
-        MIGRAPHX_DPP_REDUCE_ASM(x, v_##prefix##_u32);                                                  \
+    {                                                      \
-    }                                                                                              \
+        MIGRAPHX_DPP_REDUCE_ASM(x, v_##prefix##_f32);      \
-    __device__ inline void dpp_reduce(uint32_t& x, op) { MIGRAPHX_DPP_REDUCE_ASM(x, v_##prefix##_u32); } \
+    }                                                      \
-    __device__ inline void dpp_reduce(vec<half, 2>& x, op) { MIGRAPHX_DPP_REDUCE_ASM(x, v_pk_##prefix##_f16); }
+    __device__ inline void dpp_reduce(half& x, op)         \
+    {                                                      \
+        MIGRAPHX_DPP_REDUCE_ASM(x, v_##prefix##_f16);      \
+    }                                                      \
+    __device__ inline void dpp_reduce(int32_t& x, op)      \
+    {                                                      \
+        MIGRAPHX_DPP_REDUCE_ASM(x, v_##prefix##_u32);      \
+    }                                                      \
+    __device__ inline void dpp_reduce(uint32_t& x, op)     \
+    {                                                      \
+        MIGRAPHX_DPP_REDUCE_ASM(x, v_##prefix##_u32);      \
+    }                                                      \
+    __device__ inline void dpp_reduce(vec<half, 2>& x, op) \
+    {                                                      \
+        MIGRAPHX_DPP_REDUCE_ASM(x, v_pk_##prefix##_f16);   \
+    }
 MIGRAPHX_DPP_REDUCE(op::sum, add)
 MIGRAPHX_DPP_REDUCE(op::max, max)