Add utility functions

include/ck/utility/amd_gemm_dpp.hpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+
+#ifndef CK_AMD_GEMM_DPP_HPP
+#define CK_AMD_GEMM_DPP_HPP
+
+#include "data_type.hpp"
+
+namespace ck {
+
+// src: fp16, dst: fp32
+template <index_t MPerWave, index_t NPerWave, index_t KPerWave>
+struct intrin_dpp_f32_8x8x8_f16;
+
+template <>
+struct intrin_dpp_f32_8x8x8_f16<8, 8, 8>
+{
+    template <class FloatC>
+    __device__ static void Run(const half8_t& reg_a, const half8_t& reg_b, FloatC& reg_c)
+    {
+#if defined(__gfx1030__) || defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__)
+        reg_c.template AsType<float8_t>()(Number<0>{}) = __builtin_amdgcn_update_dpp(
+            reg_a, reg_b, reg_c.template AsType<float8_t>()[Number<0>{}]);
+#else
+        ignore = reg_a;
+        ignore = reg_b;
+        ignore = reg_c;
+#endif
+    }
+};
+
+// src: int8, dst: int32
+template <index_t MPerWave, index_t NPerWave, index_t KPerWave>
+struct intrin_dpp_i32_8x8x8_i8;
+
+template <>
+struct intrin_dpp_i32_8x8x8_i8<8, 8, 8>
+{
+    template <class IntC>
+    __device__ static void Run(const int8_t& reg_a, const int8_t& reg_b, IntC& reg_c)
+    {
+#if defined(__gfx1030__) || defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__)
+        reg_c.template AsType<int32x8_t>()(Number<0>{}) = __builtin_amdgcn_update_dpp(
+            reg_a, reg_b, reg_c.template AsType<int32x8_t>()[Number<0>{}]);
+#else
+        ignore = reg_a;
+        ignore = reg_b;
+        ignore = reg_c;
+#endif
+    }
+};
+
+} // namespace ck
+#endif