Some hiprtc fixes

src/targets/gpu/compile_hip.cpp

@@ -176,6 +176,7 @@ compile_hip_src(const std::vector<src_file>& srcs, std::string params, const std
 {
     hiprtc_program prog(srcs);
     auto options = split_string(params, ' ');
+    options.push_back("-DMIGRAPHX_HIPRTC");
     if(enabled(MIGRAPHX_GPU_DEBUG{}))
         options.push_back("-DMIGRAPHX_DEBUG");
     if(std::none_of(options.begin(), options.end(), [](const std::string& s) {
@@ -183,7 +184,7 @@ compile_hip_src(const std::vector<src_file>& srcs, std::string params, const std
        }))
         options.push_back("-std=c++17");
     options.push_back("-fno-gpu-rdc");
-    options.push_back(" -O" + string_value_of(MIGRAPHX_GPU_OPTIMIZE{}, "3"));
+    options.push_back("-O" + string_value_of(MIGRAPHX_GPU_OPTIMIZE{}, "3"));
     options.push_back("-Wno-cuda-compat");
     options.push_back("--cuda-gpu-arch=" + arch);
     prog.compile(options);
@@ -292,6 +293,8 @@ compile_hip_src(const std::vector<src_file>& srcs, std::string params, const std
     return {compiler.compile(srcs)};
 }
 
+#endif // MIGRAPHX_USE_HIPRTC
+
 std::string enum_params(std::size_t count, std::string param)
 {
     std::vector<std::string> items(count);
@@ -299,8 +302,6 @@ std::string enum_params(std::size_t count, std::string param)
     return join_strings(items, ",");
 }
 
-#endif // MIGRAPHX_USE_HIPRTC
-
 } // namespace gpu
 } // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/targets/gpu/compile_hip_code_object.cpp

@@ -80,6 +80,7 @@ std::string generate_args_hpp(const std::vector<shape>& inputs)
 
 #include <migraphx/kernels/args.hpp>
 #include <migraphx/kernels/tensor_view.hpp>
+#include <migraphx/kernels/types.hpp>
 
 namespace migraphx {
 

src/targets/gpu/kernels/include/migraphx/kernels/array.hpp

@@ -105,7 +105,7 @@ constexpr auto array_for_each(T& x, Ts&... xs)
             }
             else
             {
-                using vec_type = std::remove_reference_t<decltype(array2vec(x))>;
+                using vec_type = remove_reference_t<decltype(array2vec(x))>;
                 f(array2vec(x), __builtin_convertvector(array2vec(xs), vec_type)...);
             }
         }

src/targets/gpu/kernels/include/migraphx/kernels/hip.hpp

@@ -24,11 +24,17 @@
 #ifndef MIGRAPHX_GUARD_KERNELS_HIP_HPP
 #define MIGRAPHX_GUARD_KERNELS_HIP_HPP
 
+
+#ifndef MIGRAPHX_HIPRTC
 // Workaround macro redefinition issue with clang tidy
 #if defined(__HIP_PLATFORM_HCC__) && defined(MIGRAPHX_USE_CLANG_TIDY)
 #undef __HIP_PLATFORM_HCC__ // NOLINT
 #endif
 
 #include <hip/hip_runtime.h>
+#include <hip/hip_fp16.h>
+#include <hip/math_functions.h>
+
+#endif // MIGRAPHX_HIPRTC
 
 #endif // MIGRAPHX_GUARD_KERNELS_HIP_HPP

src/targets/gpu/kernels/include/migraphx/kernels/math.hpp

@@ -28,8 +28,7 @@
 #include <migraphx/kernels/vec.hpp>
 #include <migraphx/kernels/functional.hpp>
 #include <migraphx/kernels/type_traits.hpp>
-#include <hip/hip_fp16.h>
-#include <hip/math_functions.h>
+#include <migraphx/kernels/hip.hpp>
 
 namespace migraphx {
 

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

@@ -25,11 +25,12 @@
 #define MIGRAPHX_GUARD_AMDMIGRAPHX_KERNELS_TYPES_HPP
 
 #include <migraphx/kernels/hip.hpp>
+#include <initializer_list>
 
 namespace migraphx {
 
-using index_int = std::uint32_t;
-using diff_int  = std::int32_t;
+using index_int = uint32_t;
+using diff_int  = int32_t;
 
 #define MIGRAPHX_DEVICE_CONSTEXPR constexpr __device__ __host__ // NOLINT