fixed amd build

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

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;

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 <class T, T X, T Y>

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 {