diff --git a/composable_kernel/include/utility/common_header.hpp b/composable_kernel/include/utility/common_header.hpp index 339666ff91..b70b61a7a0 100644 --- a/composable_kernel/include/utility/common_header.hpp +++ b/composable_kernel/include/utility/common_header.hpp @@ -2,9 +2,9 @@ #define CK_COMMON_HEADER_HPP #include "config.hpp" +#include "utility.hpp" #include "integral_constant.hpp" #include "math.hpp" -#include "utility.hpp" #include "vector_type.hpp" #include "Sequence.hpp" #include "Array.hpp" @@ -12,4 +12,8 @@ #include "functional2.hpp" #include "functional3.hpp" +#if CK_USE_AMD_INLINE_ASM +#include "amd_inline_asm.hpp" +#endif + #endif diff --git a/composable_kernel/include/utility/config_amd.hpp.in b/composable_kernel/include/utility/config_amd.hpp.in index f9ae2e7830..aa12140fd7 100644 --- a/composable_kernel/include/utility/config_amd.hpp.in +++ b/composable_kernel/include/utility/config_amd.hpp.in @@ -18,10 +18,6 @@ typedef float float4_t __attribute__((ext_vector_type(4))); using index_t = uint32_t; -__device__ index_t get_thread_local_1d_id() { return threadIdx.x; } - -__device__ index_t get_block_1d_id() { return blockIdx.x; } - __device__ void fused_multiply_accumulate(float& d, const float& s0, const float& s1) { d += s0 * s1; diff --git a/composable_kernel/include/utility/integral_constant.hpp b/composable_kernel/include/utility/integral_constant.hpp index b83d7843d5..e4c213d7af 100644 --- a/composable_kernel/include/utility/integral_constant.hpp +++ b/composable_kernel/include/utility/integral_constant.hpp @@ -8,12 +8,9 @@ struct integral_constant { static constexpr T value = v; typedef T value_type; - typedef integral_constant type; // using injected-class-name + typedef integral_constant type; __host__ __device__ constexpr operator value_type() const noexcept { return value; } - __host__ __device__ constexpr value_type operator()() const noexcept - { - return value; - } // since c++14 + __host__ __device__ constexpr value_type operator()() const noexcept { return value; } }; template diff --git a/composable_kernel/include/utility/math.hpp b/composable_kernel/include/utility/math.hpp index b754ca173c..92a802a1c9 100644 --- a/composable_kernel/include/utility/math.hpp +++ b/composable_kernel/include/utility/math.hpp @@ -2,6 +2,7 @@ #define CK_MATH_HPP #include "config.hpp" +#include "integral_constant.hpp" namespace ck { namespace math {