Format

20cdddac · Paul · e9b56168 · 20cdddac
Commit 20cdddac authored Nov 21, 2023 by Paul
Show whitespace changes
Inline Side-by-side

Showing with 21 additions and 9 deletions

src/targets/gpu/kernels/include/migraphx/kernels/reduce.hpp src/targets/gpu/kernels/include/migraphx/kernels/reduce.hpp +21 -9

No files found.
--- a/src/targets/gpu/kernels/include/migraphx/kernels/reduce.hpp
+++ b/src/targets/gpu/kernels/include/migraphx/kernels/reduce.hpp
@@ -86,14 +86,26 @@ __device__ void dpp_reduce(T& in, Op op)
 // NOLINTNEXTLINE
 #define MIGRAPHX_DPP_REDUCE(op, prefix, sign)            \
-    __device__ inline void dpp_reduce(double& x, op f) { MIGRAPHX_DPP_REDUCE_ASM(x, prefix##_f64, f); } \
+    __device__ inline void dpp_reduce(double& x, op f)   \
-    __device__ inline void dpp_reduce(float& x, op f) { MIGRAPHX_DPP_REDUCE_ASM(x, prefix##_f32, f); }  \
+    {                                                    \
-    __device__ inline void dpp_reduce(half& x, op f) { MIGRAPHX_DPP_REDUCE_ASM(x, prefix##_f16, f); }   \
+        MIGRAPHX_DPP_REDUCE_ASM(x, prefix##_f64, f);     \
+    }                                                    \
+    __device__ inline void dpp_reduce(float& x, op f)    \
+    {                                                    \
+        MIGRAPHX_DPP_REDUCE_ASM(x, prefix##_f32, f);     \
+    }                                                    \
+    __device__ inline void dpp_reduce(half& x, op f)     \
+    {                                                    \
+        MIGRAPHX_DPP_REDUCE_ASM(x, prefix##_f16, f);     \
+    }                                                    \
    __device__ inline void dpp_reduce(int32_t& x, op f)  \
    {                                                    \
        MIGRAPHX_DPP_REDUCE_ASM(x, prefix##sign##32, f); \
    }                                                    \
-    __device__ inline void dpp_reduce(uint32_t& x, op f) { MIGRAPHX_DPP_REDUCE_ASM(x, prefix##_u32, f); }
+    __device__ inline void dpp_reduce(uint32_t& x, op f) \
+    {                                                    \
+        MIGRAPHX_DPP_REDUCE_ASM(x, prefix##_u32, f);     \
+    }
 // Note: when max and min are in int32_t, signed version of instruction needs to be used.
 MIGRAPHX_DPP_REDUCE(op::sum, v_add, _u)