Fix build failure for HIP-Clang (#509)

src/targets/gpu/CMakeLists.txt

@@ -69,7 +69,7 @@ add_library(migraphx_device
 set_target_properties(migraphx_device PROPERTIES EXPORT_NAME device)
 rocm_set_soversion(migraphx_device ${MIGRAPHX_SO_VERSION})
 rocm_clang_tidy_check(migraphx_device)
-target_compile_options(migraphx_device PRIVATE -std=c++17)
+target_compile_options(migraphx_device PRIVATE -std=c++17 -Xclang -fallow-half-arguments-and-returns)
 target_link_libraries(migraphx_device migraphx hip::device -Wno-invalid-command-line-argument -amdgpu-target=gfx803 -amdgpu-target=gfx900 -amdgpu-target=gfx906)
 target_include_directories(migraphx_device PUBLIC $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>)
 target_include_directories(migraphx_device PRIVATE $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/device/include>)

src/targets/gpu/device/include/migraphx/gpu/device/multi_index.hpp

@@ -31,7 +31,7 @@ struct multi_index
 };
 
 template <class ForStride>
-auto deduce_for_stride(ForStride fs) -> decltype(fs(id{}));
+__device__ __host__ auto deduce_for_stride(ForStride fs) -> decltype(fs(id{}));
 
 MIGRAPHX_DEVICE_CONSTEXPR multi_index<1> make_multi_index(index_int i, index_int n)
 {

src/targets/gpu/device/include/migraphx/gpu/device/reduce.hpp

@@ -69,7 +69,7 @@ struct min
 struct lowest
 {
     template <class T>
-    operator T() const
+    __device__ __host__ operator T() const
     {
         return device_cast(std::numeric_limits<host_type<T>>::lowest());
     }
@@ -78,7 +78,7 @@ struct lowest
 struct highest
 {
     template <class T>
-    operator T() const
+    __device__ __host__ operator T() const
     {
         return device_cast(std::numeric_limits<host_type<T>>::max());
     }

src/targets/gpu/device/include/migraphx/gpu/device/types.hpp

@@ -103,19 +103,19 @@ host_type<T>* host_cast(T* x)
 }
 
 template <class T>
-device_type<T> device_cast(const T& x)
+__device__ __host__ device_type<T> device_cast(const T& x)
 {
     return reinterpret_cast<const device_type<T>&>(x);
 }
 
 template <class T>
-device_type<T>* device_cast(T* x)
+__device__ __host__ device_type<T>* device_cast(T* x)
 {
     return reinterpret_cast<device_type<T>*>(x);
 }
 
 template <class T>
-tensor_view<device_type<T>> device_cast(tensor_view<T> x)
+__device__ __host__ tensor_view<device_type<T>> device_cast(tensor_view<T> x)
 {
     return {x.get_shape(), reinterpret_cast<device_type<T>*>(x.data())};
 }