SWDEV-395996 - Add HIPRTC support for missing hip headers

hip headers such as hip_math_constants.h, library_types.h,
hip_common.h, channel_descriptor.h, device_functions.h, hip_complex.h,
hip_texture_types.h, math_functions.h, surface_types.h
are added in HIPRTC

Change-Id: I4a4c198449ceb609c3ff55e00b43056c1f085431


[ROCm/clr commit: d7d0f11318]

projects/clr/hipamd/CMakeLists.txt

@@ -292,15 +292,15 @@ set(_versionInfoHeader
 ")
 file(WRITE "${PROJECT_BINARY_DIR}/include/hip/hip_version.h" ${_versionInfoHeader})
 
-if(HIP_RUNTIME STREQUAL "rocclr")
-	add_subdirectory(src)
-endif()
-
 # Generate .hipInfo
-file(WRITE "${PROJECT_BINARY_DIR}/.hipInfo" ${_buildInfo})
+file(WRITE "${PROJECT_BINARY_DIR}/share/hip/.hipInfo" ${_buildInfo})
 
 # Generate version
-file(WRITE "${PROJECT_BINARY_DIR}/version" ${_versionInfo})
+file(WRITE "${PROJECT_BINARY_DIR}/share/hip/version" ${_versionInfo})
+
+if(HIP_RUNTIME STREQUAL "rocclr")
+   add_subdirectory(src)
+endif()
 
 # Build doxygen documentation
 find_program(DOXYGEN_EXE doxygen)
@@ -321,14 +321,14 @@ endif()
 #############################
 
 # Install .hipInfo
-install(FILES ${PROJECT_BINARY_DIR}/.hipInfo DESTINATION ${CMAKE_INSTALL_LIBDIR})
+install(FILES ${PROJECT_BINARY_DIR}/share/hip/.hipInfo DESTINATION ${CMAKE_INSTALL_LIBDIR})
 
 # Install version
-install(FILES ${PROJECT_BINARY_DIR}/version DESTINATION ${CMAKE_INSTALL_DATADIR}/hip)
+install(FILES ${PROJECT_BINARY_DIR}/share/hip/version DESTINATION ${CMAKE_INSTALL_DATADIR}/hip)
 # .hipVersion is added to satisfy Windows compute build.
 #TODO to be removed
 if(WIN32)
-    install(FILES ${PROJECT_BINARY_DIR}/version DESTINATION ${CMAKE_INSTALL_BINDIR} RENAME .hipVersion)
+    install(FILES ${PROJECT_BINARY_DIR}/share/hip/version DESTINATION ${CMAKE_INSTALL_BINDIR} RENAME .hipVersion)
 endif()
 
 # Install src, bin, include & cmake if necessary

projects/clr/hipamd/include/hip/amd_detail/amd_channel_descriptor.h

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015 - 2022 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -23,9 +23,16 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_CHANNEL_DESCRIPTOR_H
 #define HIP_INCLUDE_HIP_AMD_DETAIL_CHANNEL_DESCRIPTOR_H
 
+#if defined(__clang__)
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wold-style-cast"
+#endif
+
+#if !defined(__HIPCC_RTC__)
 #include <hip/hip_common.h>
 #include <hip/driver_types.h>
 #include <hip/amd_detail/amd_hip_vector_types.h>
+#endif
 
 #ifdef __cplusplus
 
@@ -353,4 +360,8 @@ struct hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w,
 
 #endif /* __cplusplus */
 
+#if defined(__clang__)
+#pragma clang diagnostic pop
+#endif
+
 #endif /* !HIP_INCLUDE_HIP_AMD_DETAIL_CHANNEL_DESCRIPTOR_H */

projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h

@@ -23,18 +23,32 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_DEVICE_FUNCTIONS_H
 #define HIP_INCLUDE_HIP_AMD_DETAIL_DEVICE_FUNCTIONS_H
 
-#include "host_defines.h"
-#include "math_fwd.h"
+#if defined(__clang__)
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wreserved-macro-identifier"
+#pragma clang diagnostic ignored "-Wreserved-identifier"
+#pragma clang diagnostic ignored "-Wsign-conversion"
+#pragma clang diagnostic ignored "-Wc++98-compat-pedantic"
+#pragma clang diagnostic ignored "-Wold-style-cast"
+#pragma clang diagnostic ignored "-Wshorten-64-to-32"
+#pragma clang diagnostic ignored "-Wimplicit-int-conversion"
+#pragma clang diagnostic ignored "-Wimplicit-float-conversion"
+#pragma clang diagnostic ignored "-Wmissing-noreturn"
+#pragma clang diagnostic ignored "-Wimplicit-fallthrough"
+#pragma clang diagnostic ignored "-Wunneeded-internal-declaration"
+#pragma clang diagnostic ignored "-Wshift-count-overflow"
+#endif
 
 #if !defined(__HIPCC_RTC__)
+#include "host_defines.h"
+#include "math_fwd.h"
 #include <hip/hip_runtime_api.h>
 #include <stddef.h>
-#endif // !defined(__HIPCC_RTC__)
-
 #include <hip/hip_vector_types.h>
 #include <hip/amd_detail/device_library_decls.h>
+#endif // !defined(__HIPCC_RTC__)
 
-#if __HIP_CLANG_ONLY__
+#if defined(__clang__) && defined(__HIP__)
 extern "C" __device__ int printf(const char *fmt, ...);
 #else
 template <typename... All>
@@ -639,7 +653,7 @@ __device__ static inline float __ull2float_rz(unsigned long long int x) {
     return __ocml_cvtrtz_f32_u64(x);
 }
 
-#if __HIP_CLANG_ONLY__
+#if defined(__clang__) && defined(__HIP__)
 
 // Clock functions
 __device__ long long int __clock64();
@@ -1105,4 +1119,9 @@ static inline __device__ void* memset(void* ptr, int val, size_t size) {
     return __hip_hc_memset(ptr, val8, size);
 }
 #endif // !__OPENMP_AMDGCN__
+
+#if defined(__clang__)
+#pragma clang diagnostic pop
+#endif
+
 #endif

projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h

@@ -22,7 +22,17 @@ THE SOFTWARE.
 
 #pragma once
 
+#if defined(__clang__)
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wc++98-compat"
+#pragma clang diagnostic ignored "-Wold-style-cast"
+#pragma clang diagnostic ignored "-Wc++98-compat-pedantic"
+#pragma clang diagnostic ignored "-Wsign-conversion"
+#endif
+
+#if !defined(__HIPCC_RTC__)
 #include "amd_device_functions.h"
+#endif
 
 #if __has_builtin(__hip_atomic_compare_exchange_strong)
 
@@ -1625,4 +1635,8 @@ unsigned long long atomicXor(
     return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED);
 }
 
+#if defined(__clang__)
+#pragma clang diagnostic pop
+#endif
+
 #endif // __hip_atomic_compare_exchange_strong

projects/clr/hipamd/include/hip/amd_detail/amd_hip_complex.h

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -23,7 +23,21 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
 #define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
 
+#if defined(__clang__)
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wfloat-equal"
+#pragma clang diagnostic ignored "-Wc++98-compat"
+#pragma clang diagnostic ignored "-Wold-style-cast"
+#pragma clang diagnostic ignored "-Wreserved-macro-identifier"
+#pragma clang diagnostic ignored "-Wimplicit-int-float-conversion"
+#pragma clang diagnostic ignored "-Wimplicit-float-conversion"
+#pragma clang diagnostic ignored "-Wdouble-promotion"
+#pragma clang diagnostic ignored "-Wc++98-compat-pedantic"
+#endif
+
+#if !defined(__HIPCC_RTC__)
 #include "hip/amd_detail/amd_hip_vector_types.h"
+#endif
 
 #if defined(__HIPCC_RTC__)
 #define __HOST_DEVICE__ __device__
@@ -311,4 +325,8 @@ __HOST_DEVICE__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDo
     return make_hipDoubleComplex(real, imag);
 }
 
+#if defined(__clang__)
+#pragma clang diagnostic pop
+#endif
+
 #endif //HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H

projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime.h

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -110,18 +110,19 @@ using ::int64_t;
 
 #if !defined(__HIPCC_RTC__)
 #include <hip/hip_runtime_api.h>
+#include <hip/amd_detail/amd_hip_atomic.h>
+#include <hip/amd_detail/amd_device_functions.h>
+#include <hip/amd_detail/amd_surface_functions.h>
+#include <hip/amd_detail/texture_fetch_functions.h>
+#include <hip/amd_detail/texture_indirect_functions.h>
 extern int HIP_TRACE_API;
 #endif // !defined(__HIPCC_RTC__)
 
 #ifdef __cplusplus
 #include <hip/amd_detail/hip_ldg.h>
 #endif
-#include <hip/amd_detail/amd_hip_atomic.h>
+
 #include <hip/amd_detail/host_defines.h>
-#include <hip/amd_detail/amd_device_functions.h>
-#include <hip/amd_detail/amd_surface_functions.h>
-#include <hip/amd_detail/texture_fetch_functions.h>
-#include <hip/amd_detail/texture_indirect_functions.h>
 
 // TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
 #if defined(__KALMAR_ACCELERATOR__) && !defined(__HCC_ACCELERATOR__)
@@ -244,10 +245,10 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
 #include <hip/hip_runtime_api.h>
 #endif // !defined(__HIPCC_RTC__)
 
-extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_id(uint);
-extern "C" __device__ __attribute__((const)) size_t __ockl_get_group_id(uint);
-extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_size(uint);
-extern "C" __device__ __attribute__((const)) size_t __ockl_get_num_groups(uint);
+extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_id(unsigned int);
+extern "C" __device__ __attribute__((const)) size_t __ockl_get_group_id(unsigned int);
+extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_size(unsigned int);
+extern "C" __device__ __attribute__((const)) size_t __ockl_get_num_groups(unsigned int);
 struct __HIP_BlockIdx {
   __device__
   std::uint32_t operator()(std::uint32_t x) const noexcept { return __ockl_get_group_id(x); }
@@ -282,7 +283,7 @@ typedef struct dim3 {
 } dim3;
 #endif // !defined(__HIPCC_RTC__)
 
-extern "C" __device__ __attribute__((const)) size_t __ockl_get_global_size(uint);
+extern "C" __device__ __attribute__((const)) size_t __ockl_get_global_size(unsigned int);
 
 #ifdef __cplusplus
 template <typename F> struct __HIP_Coordinates {
@@ -367,33 +368,35 @@ static constexpr __HIP_Coordinates<__HIP_GridDim> gridDim{};
 static constexpr __HIP_Coordinates<__HIP_ThreadIdx> threadIdx{};
 #endif // __cplusplus
 
-extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_id(uint);
+extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_id(unsigned int);
 #define hipThreadIdx_x (__ockl_get_local_id(0))
 #define hipThreadIdx_y (__ockl_get_local_id(1))
 #define hipThreadIdx_z (__ockl_get_local_id(2))
 
-extern "C" __device__ __attribute__((const)) size_t __ockl_get_group_id(uint);
+extern "C" __device__ __attribute__((const)) size_t __ockl_get_group_id(unsigned int);
 #define hipBlockIdx_x (__ockl_get_group_id(0))
 #define hipBlockIdx_y (__ockl_get_group_id(1))
 #define hipBlockIdx_z (__ockl_get_group_id(2))
 
-extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_size(uint);
+extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_size(unsigned int);
 #define hipBlockDim_x (__ockl_get_local_size(0))
 #define hipBlockDim_y (__ockl_get_local_size(1))
 #define hipBlockDim_z (__ockl_get_local_size(2))
 
-extern "C" __device__ __attribute__((const)) size_t __ockl_get_num_groups(uint);
+extern "C" __device__ __attribute__((const)) size_t __ockl_get_num_groups(unsigned int);
 #define hipGridDim_x (__ockl_get_num_groups(0))
 #define hipGridDim_y (__ockl_get_num_groups(1))
 #define hipGridDim_z (__ockl_get_num_groups(2))
 
+#if !defined(__HIPCC_RTC__)
 #include <hip/amd_detail/amd_math_functions.h>
+#endif
 
 #if __HIP_HCC_COMPAT_MODE__
 // Define HCC work item functions in terms of HIP builtin variables.
 #pragma push_macro("__DEFINE_HCC_FUNC")
 #define __DEFINE_HCC_FUNC(hc_fun,hip_var) \
-inline __device__ __attribute__((always_inline)) uint hc_get_##hc_fun(uint i) { \
+inline __device__ __attribute__((always_inline)) unsigned int hc_get_##hc_fun(unsigned int i) { \
   if (i==0) \
     return hip_var.x; \
   else if(i==1) \
@@ -408,11 +411,11 @@ __DEFINE_HCC_FUNC(group_size, blockDim)
 __DEFINE_HCC_FUNC(num_groups, gridDim)
 #pragma pop_macro("__DEFINE_HCC_FUNC")
 
-extern "C" __device__ __attribute__((const)) size_t __ockl_get_global_id(uint);
-inline __device__ __attribute__((always_inline)) uint
+extern "C" __device__ __attribute__((const)) size_t __ockl_get_global_id(unsigned int);
+inline __device__ __attribute__((always_inline)) unsigned int
 hc_get_workitem_absolute_id(int dim)
 {
-  return (uint)__ockl_get_global_id(dim);
+  return (unsigned int)__ockl_get_global_id(dim);
 }
 
 #endif

projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2018 - 2022 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2018 - 2023 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -23,12 +23,33 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_SURFACE_FUNCTIONS_H
 #define HIP_INCLUDE_HIP_AMD_DETAIL_SURFACE_FUNCTIONS_H
 
+#if defined(__clang__)
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wreserved-macro-identifier"
+#pragma clang diagnostic ignored "-Wc++17-extensions"
+#pragma clang diagnostic ignored "-Wreserved-identifier"
+#pragma clang diagnostic ignored "-Wc++98-compat"
+#pragma clang diagnostic ignored "-Wold-style-cast"
+#pragma clang diagnostic ignored "-Wextra-semi-stmt"
+#pragma clang diagnostic ignored "-Wunused-variable"
+#pragma clang diagnostic ignored "-Wunused-parameter"
+#pragma clang diagnostic ignored "-Wunused-template"
+#endif
+
 #if defined(__cplusplus)
 
+#if !defined(__HIPCC_RTC__)
 #include <hip/surface_types.h>
 #include <hip/hip_vector_types.h>
 #include <hip/amd_detail/texture_fetch_functions.h>
 #include <hip/amd_detail/ockl_image.h>
+#endif
+
+#if defined(__HIPCC_RTC__)
+#define __HOST_DEVICE__ __device__
+#else
+#define __HOST_DEVICE__ __host__ __device__
+#endif
 
 #define __HIP_SURFACE_OBJECT_PARAMETERS_INIT                                                            \
     unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)surfObj; 
@@ -232,4 +253,9 @@ static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfa
 }
 
 #endif
+
+#if defined(__clang__)
+#pragma clang diagnostic pop
+#endif
+
 #endif

projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -31,7 +31,16 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_DEVICE_LIBRARY_DECLS_H
 #define HIP_INCLUDE_HIP_AMD_DETAIL_DEVICE_LIBRARY_DECLS_H
 
+#if defined(__clang__)
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wc++98-compat-pedantic"
+#pragma clang diagnostic ignored "-Wreserved-identifier"
+#pragma clang diagnostic ignored "-Wold-style-cast"
+#endif
+
+#if !defined(__HIPCC_RTC__)
 #include "hip/amd_detail/host_defines.h"
+#endif
 
 typedef unsigned char uchar;
 typedef unsigned short ushort;
@@ -128,4 +137,8 @@ __device__ inline static __local void* __to_local(unsigned x) { return (__local
 #define __CLK_LOCAL_MEM_FENCE    0x01
 typedef unsigned __cl_mem_fence_flags;
 
+#if defined(__clang__)
+#pragma clang diagnostic pop
+#endif
+
 #endif

projects/clr/hipamd/include/hip/amd_detail/math_fwd.h

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -22,13 +22,22 @@ THE SOFTWARE.
 
 #pragma once
 
+#if defined(__clang__)
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wundef"
+#pragma clang diagnostic ignored "-Wreserved-identifier"
+#endif
+
+#if !defined(__HIPCC_RTC__)
 #include "host_defines.h"
+#endif
+
 #if defined(__cplusplus)
     extern "C" {
 #endif
 
 // DOT FUNCTIONS
-#if __HIP_CLANG_ONLY__
+#if defined(__clang__) && defined(__HIP__)
 __device__
 __attribute__((const))
 int __ockl_sdot2(
@@ -692,3 +701,7 @@ double __ocml_fma_rtz_f64(double, double, double);
 #if defined(__cplusplus)
     } // extern "C"
 #endif
+
+#if defined(__clang__)
+#pragma clang diagnostic pop
+#endif

projects/clr/hipamd/include/hip/amd_detail/ockl_image.h

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015 - 2022 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -22,7 +22,14 @@ THE SOFTWARE.
 
 #pragma once
 
+#if defined(__clang__)
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wreserved-identifier"
+#endif
+
+#if !defined(__HIPCC_RTC__)
 #include <hip/hip_vector_types.h>
+#endif
 
 extern "C" {
 
@@ -172,4 +179,8 @@ __device__ int __ockl_image_channel_order_CM(unsigned int ADDRESS_SPACE_CONSTANT
 
 __device__ int __ockl_image_channel_order_CMa(unsigned int ADDRESS_SPACE_CONSTANT* i);
 
-};
+}
+
+#if defined(__clang__)
+#pragma clang diagnostic pop
+#endif

projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -22,13 +22,24 @@ THE SOFTWARE.
 
 #pragma once
 
+#if defined(__clang__)
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wc++17-extensions"
+#pragma clang diagnostic ignored "-Wreserved-identifier"
+#pragma clang diagnostic ignored "-Wc++98-compat"
+#pragma clang diagnostic ignored "-Wold-style-cast"
+#pragma clang diagnostic ignored "-Wextra-semi-stmt"
+#pragma clang diagnostic ignored "-Wunused-variable"
+#pragma clang diagnostic ignored "-Wunused-parameter"
+#pragma clang diagnostic ignored "-Wunused-template"
+#endif
+
 #if defined(__cplusplus)
 
+#if !defined(__HIPCC_RTC__)
 #include <hip/hip_vector_types.h>
 #include <hip/hip_texture_types.h>
 #include <hip/amd_detail/ockl_image.h>
-
-#if !defined(__HIPCC_RTC__)
 #include <type_traits>
 #endif // !defined(__HIPCC_RTC__)
 
@@ -490,3 +501,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex2dgather_ret_t<T, rea
 }
 
 #endif
+
+#if defined(__clang__)
+#pragma clang diagnostic pop
+#endif

projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -22,14 +22,24 @@ THE SOFTWARE.
 
 #pragma once
 
+#if defined(__clang__)
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wc++17-extensions"
+#pragma clang diagnostic ignored "-Wreserved-identifier"
+#pragma clang diagnostic ignored "-Wc++98-compat"
+#pragma clang diagnostic ignored "-Wold-style-cast"
+#pragma clang diagnostic ignored "-Wunused-variable"
+#pragma clang diagnostic ignored "-Wunused-parameter"
+#pragma clang diagnostic ignored "-Wunused-template"
+#endif
+
 #if defined(__cplusplus)
 
+#if !defined(__HIPCC_RTC__)
 #include <hip/hip_vector_types.h>
 #include <hip/hip_texture_types.h>
 #include <hip/amd_detail/texture_fetch_functions.h>
 #include <hip/amd_detail/ockl_image.h>
-
-#if !defined(__HIPCC_RTC__)
 #include <type_traits>
 #endif // !defined(__HIPCC_RTC__)
 
@@ -208,7 +218,7 @@ static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject
         return __hipMapFrom<T>(tmp);
         break;
     }
-    };
+    }
     return {};
 }
 
@@ -477,3 +487,7 @@ static __device__ __hip_img_chk__ void texCubemapLayeredGrad(T *ptr, hipTextureO
 }
 
 #endif
+
+#if defined(__clang__)
+#pragma clang diagnostic pop
+#endif

projects/clr/hipamd/packaging/CMakeLists.txt

@@ -60,7 +60,7 @@ if(HIP_PLATFORM STREQUAL "amd" )
 
 #TODO:This do not belong in BINARY package.
 #Keeping it as is for now
-install(FILES ${CMAKE_BINARY_DIR}/hipamd/.hipInfo DESTINATION ${CMAKE_INSTALL_LIBDIR} COMPONENT binary)
+install(FILES ${CMAKE_BINARY_DIR}/hipamd/share/hip/.hipInfo DESTINATION ${CMAKE_INSTALL_LIBDIR} COMPONENT binary)
 
 install ( EXPORT hip-targets  FILE hip-targets.cmake NAMESPACE hip:: DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/hip  COMPONENT binary)
 install ( EXPORT hip-lang-targets  FILE hip-lang-targets.cmake NAMESPACE hip-lang:: DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/hip-lang  COMPONENT binary)
@@ -94,11 +94,11 @@ install(FILES ${CMAKE_BINARY_DIR}/hipamd/include/hip/amd_detail/hip_prof_str.h
         DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/hip/amd_detail COMPONENT dev)
 install(FILES ${CMAKE_BINARY_DIR}/hipamd/include/hip/hip_version.h
         DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/hip COMPONENT dev)
-install(FILES ${CMAKE_BINARY_DIR}/hipamd/version DESTINATION ${CMAKE_INSTALL_DATADIR}/hip COMPONENT dev)
+install(FILES ${CMAKE_BINARY_DIR}/hipamd/share/hip/version DESTINATION ${CMAKE_INSTALL_DATADIR}/hip COMPONENT dev)
 # .hipVersion is added to satisfy Windows compute build.
 #TODO to be removed
 if(WIN32)
-    install(FILES ${CMAKE_BINARY_DIR}/hipamd/version DESTINATION ${CMAKE_INSTALL_BINDIR} RENAME .hipVersion COMPONENT dev)
+    install(FILES ${CMAKE_BINARY_DIR}/hipamd/share/hip/version DESTINATION ${CMAKE_INSTALL_BINDIR} RENAME .hipVersion COMPONENT dev)
 endif()
 install(DIRECTORY ${HIP_COMMON_DIR}/cmake/ DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/hip COMPONENT dev)
 install(FILES ${CMAKE_BINARY_DIR}/hipamd/hip-config.cmake ${CMAKE_BINARY_DIR}/hipamd/hip-config-version.cmake DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/hip COMPONENT dev)

projects/clr/hipamd/src/CMakeLists.txt

@@ -261,9 +261,11 @@ endif()
 target_compile_definitions(amdhip64 PUBLIC USE_PROF_API=1)
 
 add_custom_command(TARGET amdhip64 POST_BUILD COMMAND
-  ${CMAKE_COMMAND} -E copy ${PROJECT_BINARY_DIR}/.hipInfo ${PROJECT_BINARY_DIR}/lib/.hipInfo)
+  ${CMAKE_COMMAND} -E copy ${PROJECT_BINARY_DIR}/share/hip/.hipInfo ${PROJECT_BINARY_DIR}/lib/.hipInfo)
 add_custom_command(TARGET amdhip64 POST_BUILD COMMAND
   ${CMAKE_COMMAND} -E copy_directory ${PROJECT_SOURCE_DIR}/include ${PROJECT_BINARY_DIR}/include)
+add_custom_command(TARGET amdhip64 POST_BUILD COMMAND
+  ${CMAKE_COMMAND} -E copy_directory ${HIP_COMMON_INCLUDE_DIR} ${PROJECT_BINARY_DIR}/include)
 
 add_library(host INTERFACE)
 target_link_libraries(host INTERFACE amdhip64)

projects/clr/hipamd/src/hiprtc/CMakeLists.txt

@@ -1,4 +1,4 @@
-# Copyright (c) 2020 - 2022 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2020 - 2023 Advanced Micro Devices, Inc. All rights reserved.
 #
 # Permission is hereby granted, free of charge, to any person obtaining a copy
 # of this software and associated documentation files (the "Software"), to deal
@@ -129,12 +129,33 @@ set(HIPRTC_GEN_HEADER "${HIPRTC_GEN_DIR}/hipRTC_header.h")
 set(HIPRTC_GEN_MCIN "${HIPRTC_GEN_DIR}/hipRTC_header.mcin")
 set(HIPRTC_GEN_PREPROCESSED "${HIPRTC_GEN_DIR}/hipRTC")
 set(HIPRTC_GEN_OBJ "${HIPRTC_GEN_DIR}/hipRTC_header${CMAKE_CXX_OUTPUT_EXTENSION}")
-set(HIPRTC_WARP_FUNCS "${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_warp_functions.h")
-set(HIPRTC_FP16_MATH_FWD "${PROJECT_SOURCE_DIR}/include/hip/amd_detail/hip_fp16_math_fwd.h")
-set(HIPRTC_FP16_FUNCS "${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_fp16.h")
-set(HIPRTC_COOP_GROUPS "${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_cooperative_groups.h")
-set(HIPRTC_COOP_GRPS_HELPER "${PROJECT_SOURCE_DIR}/include/hip/amd_detail/hip_cooperative_groups_helper.h")
-set(HIPRTC_UNSAFE_ATOMICS "${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_unsafe_atomics.h")
+# list of headers which needs to be appended to the hiprtc preprocessed file
+set(HIPRTC_HEADERS
+${HIP_COMMON_INCLUDE_DIR}/hip/hip_common.h
+${HIP_COMMON_INCLUDE_DIR}/hip/library_types.h
+${HIP_COMMON_INCLUDE_DIR}/hip/driver_types.h
+${HIP_COMMON_INCLUDE_DIR}/hip/surface_types.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_channel_descriptor.h
+${HIP_COMMON_INCLUDE_DIR}/hip/texture_types.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/ockl_image.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/texture_fetch_functions.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/texture_indirect_functions.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_surface_functions.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_complex.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_math_constants.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/math_fwd.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/device_library_decls.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_device_functions.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_warp_functions.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/hip_cooperative_groups_helper.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_cooperative_groups.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_unsafe_atomics.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_atomic.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/math_fwd.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/hip_fp16_math_fwd.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_fp16.h
+${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_math_functions.h
+)
 
 # Generate required HIPRTC files.
 FILE(MAKE_DIRECTORY ${HIPRTC_GEN_DIR})
@@ -147,7 +168,7 @@ generate_hiprtc_mcin("${HIPRTC_GEN_MCIN}" "${HIPRTC_GEN_PREPROCESSED}")
 add_custom_command(
   OUTPUT ${HIPRTC_GEN_PREPROCESSED}
   COMMAND $<TARGET_FILE:clang> -O3 --rocm-path=${PROJECT_SOURCE_DIR}/include/.. -std=c++17 -nogpulib --hip-version=3.6 -isystem ${HIP_COMMON_INCLUDE_DIR} -isystem ${PROJECT_SOURCE_DIR}/include -isystem ${PROJECT_BINARY_DIR}/include -isystem ${CMAKE_CURRENT_SOURCE_DIR}/include --cuda-device-only -D__HIPCC_RTC__ -x hip ${HIPRTC_GEN_HEADER} -E -o ${HIPRTC_GEN_PREPROCESSED}
-  COMMAND ${CMAKE_COMMAND} -DHIPRTC_ADD_MACROS=1 -DHIPRTC_WARP_HEADER_FILE=${HIPRTC_WARP_FUNCS} -DHIPRTC_COOP_HEADER_FILE=${HIPRTC_COOP_GROUPS} -DHIPRTC_COOP_HELPER_FILE=${HIPRTC_COOP_GRPS_HELPER} -DHIPRTC_UNSAFE_ATOMICS_FILE=${HIPRTC_UNSAFE_ATOMICS} -DHIPRTC_FP16_MATH_FWD_FILE=${HIPRTC_FP16_MATH_FWD} -DHIPRTC_FP16_HEADER_FILE=${HIPRTC_FP16_FUNCS} -DHIPRTC_PREPROCESSED_FILE=${HIPRTC_GEN_PREPROCESSED} -P ${HIPRTC_CMAKE}
+  COMMAND ${CMAKE_COMMAND} -DHIPRTC_ADD_MACROS=1 -DHIPRTC_HEADERS="${HIPRTC_HEADERS}" -DHIPRTC_PREPROCESSED_FILE=${HIPRTC_GEN_PREPROCESSED} -P ${HIPRTC_CMAKE}
   DEPENDS clang ${HIPRTC_GEN_HEADER})
 add_custom_command(
   OUTPUT ${HIPRTC_GEN_OBJ}

projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake

@@ -30,6 +30,8 @@ function(get_hiprtc_macros HIPRTC_DEFINES)
 "#pragma clang diagnostic push\n\
 #pragma clang diagnostic ignored \"-Wreserved-id-macro\"\n\
 #pragma clang diagnostic ignored \"-Wc++98-compat-pedantic\"\n\
+#pragma clang diagnostic ignored \"-Wreserved-macro-identifier\"\n\
+#pragma clang diagnostic ignored \"-Wundef\"\n\
 #define __device__ __attribute__((device))\n\
 #define __host__ __attribute__((host))\n\
 #define __global__ __attribute__((global))\n\
@@ -40,7 +42,11 @@ function(get_hiprtc_macros HIPRTC_DEFINES)
 #define __noinline__ __attribute__((noinline))\n\
 #endif\n\
 #define __forceinline__ inline __attribute__((always_inline))\n\
-
+#if __HIP_NO_IMAGE_SUPPORT\n\
+#define __hip_img_chk__ __attribute__((unavailable(\"The image/texture API not supported on the device\")))\n\
+#else\n\
+#define __hip_img_chk__\n\
+#endif\n\
 #define launch_bounds_impl0(requiredMaxThreadsPerBlock)                                       \\\n\
     __attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock)))\n\
 #define launch_bounds_impl1(requiredMaxThreadsPerBlock, minBlocksPerMultiprocessor)           \\\n\
@@ -49,13 +55,11 @@ function(get_hiprtc_macros HIPRTC_DEFINES)
 #define select_impl_(_1, _2, impl_, ...) impl_\n\
 #define __launch_bounds__(...)                                                                \\\n\
     select_impl_(__VA_ARGS__, launch_bounds_impl1, launch_bounds_impl0)(__VA_ARGS__)           \n\
-#pragma clang diagnostic pop\n\
 #define HIP_INCLUDE_HIP_HIP_RUNTIME_H\n\
-#pragma clang diagnostic push\n\
-#pragma clang diagnostic ignored \"-Wreserved-macro-identifier\"\n\
 #define _HIP_BFLOAT16_H_\n\
-#pragma clang diagnostic pop\n\
-#define HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H"
+#define HIP_INCLUDE_HIP_MATH_FUNCTIONS_H\n\
+#define HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H\n\
+#pragma clang diagnostic pop"
   PARENT_SCOPE)
 endfunction(get_hiprtc_macros)
 
@@ -64,20 +68,15 @@ if(HIPRTC_ADD_MACROS)
   message(STATUS "Appending hiprtc macros to ${HIPRTC_PREPROCESSED_FILE}.")
   get_hiprtc_macros(HIPRTC_DEFINES)
   FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_DEFINES}")
-  FILE(READ "${HIPRTC_WARP_HEADER_FILE}" HIPRTC_WARP_HEADER)
-  FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_WARP_HEADER}")
 #pragma clang diagnostic push
 #pragma clang diagnostic ignored "-Wreserved-macro-identifier"
-  FILE(READ "${HIPRTC_COOP_HELPER_FILE}" HIPRTC_COOP_HELPER)
-  FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_COOP_HELPER}")
-  FILE(READ "${HIPRTC_COOP_HEADER_FILE}" HIPRTC_COOP_HEADER)
-  FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_COOP_HEADER}")
-  FILE(READ "${HIPRTC_UNSAFE_ATOMICS_FILE}" HIPRTC_UNSAFE_ATOMICS)
-  FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_UNSAFE_ATOMICS}")
-  FILE(READ "${HIPRTC_FP16_MATH_FWD_FILE}" HIPRTC_FP16_MATH_FWD)
-  FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_FP16_MATH_FWD}")
-  FILE(READ "${HIPRTC_FP16_HEADER_FILE}" HIPRTC_FP16_HEADER)
-  FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_FP16_HEADER}")
+  set(HIPRTC_HEADER_LIST ${HIPRTC_HEADERS})
+  separate_arguments(HIPRTC_HEADER_LIST)
+# appends all the headers from the list to the hiprtc preprocessed file
+  foreach(header ${HIPRTC_HEADER_LIST})
+    FILE(READ "${header}" HEADER_FILE)
+    FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HEADER_FILE}")
+  endforeach()
 #pragma clang diagnostic pop
 endif()