SWDEV-288707 - Merge HIP refactored code to staging

These changes move-
- rocclr based implementation inside src/hipamd/src
- platform specific (both nvidia and amd) headers inside
src/hipamd/include/hip

Change-Id: Ia29791a727244952591fe1d813dcef0303b73a9e


[ROCm/hip commit: ce76ec8b88]

projects/hip/CMakeLists.txt

@@ -285,7 +285,7 @@ set(_versionInfoHeader
 file(WRITE "${PROJECT_BINARY_DIR}/include/hip/hip_version.h" ${_versionInfoHeader})
 
 if(HIP_RUNTIME STREQUAL "rocclr")
-    add_subdirectory(rocclr)
+    add_subdirectory(src/hipamd)
 endif()
 
 # Generate .hipInfo
@@ -322,6 +322,7 @@ if(NOT ${INSTALL_SOURCE} EQUAL 0)
     install(CODE "file(REMOVE_RECURSE ${CMAKE_INSTALL_PREFIX}/include)")
 
     install(DIRECTORY include DESTINATION .)
+    install(DIRECTORY src/hipamd/include/hip/ DESTINATION include/hip/)
     install(DIRECTORY cmake DESTINATION .)
 endif()
 
@@ -478,6 +479,7 @@ set(HIP_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR})
 if(HIP_PLATFORM STREQUAL "nvidia")
     execute_process(COMMAND "${CMAKE_COMMAND}" -E copy_directory "${HIP_SRC_PATH}/include" "${HIP_ROOT_DIR}/include" RESULT_VARIABLE RUN_HIT ERROR_QUIET)
 endif()
+execute_process(COMMAND "${CMAKE_COMMAND}" -E copy_directory "${HIP_SRC_PATH}/src/hipamd/include/hip/" "${HIP_ROOT_DIR}/include/hip/" RESULT_VARIABLE RUN_HIT ERROR_QUIET)
 execute_process(COMMAND "${CMAKE_COMMAND}" -E copy_directory "${HIP_SRC_PATH}/cmake" "${HIP_ROOT_DIR}/cmake" RESULT_VARIABLE RUN_HIT ERROR_QUIET)
 if(${RUN_HIT} EQUAL 0)
     execute_process(COMMAND "${CMAKE_COMMAND}" -E copy_directory "${HIP_SRC_PATH}/bin" "${HIP_ROOT_DIR}/bin" RESULT_VARIABLE RUN_HIT ERROR_QUIET)

projects/hip/bin/hip_embed_pch.sh

@@ -21,7 +21,7 @@
 
 printUsage() {
   echo
-  echo "Usage: $(basename "$0") HIP_BUILD_INC_DIR HIP_INC_DIR LLVM_DIR [option] [RTC_LIB_OUTPUT]"
+  echo "Usage: $(basename "$0") HIP_BUILD_INC_DIR HIP_INC_DIR HIP_AMD_INC_DIR LLVM_DIR [option] [RTC_LIB_OUTPUT]"
   echo
   echo "Options:"
   echo "  -p,  --generate_pch  Generate pre-compiled header (default)"
@@ -39,13 +39,14 @@ fi
 
 HIP_BUILD_INC_DIR="$1"
 HIP_INC_DIR="$2"
-LLVM_DIR="$3"
+HIP_AMD_INC_DIR="$3"
+LLVM_DIR="$4"
 # By default, generate pch
 TARGET="generatepch"
 
-while [ "$4" != "" ];
+while [ "$5" != "" ];
 do
-  case "$4" in
+  case "$5" in
     -h | --help )
         printUsage ; exit 0 ;;
     -p | --generate_pch )
@@ -58,9 +59,9 @@ do
   shift 1
 done
 
-# Allow hiprtc lib name to be set by argument 6
-if [[ "$5" != "" ]]; then
-  rtc_shared_lib_out="$5"
+# Allow hiprtc lib name to be set by argument 7
+if [[ "$6" != "" ]]; then
+  rtc_shared_lib_out="$6"
 else
   if [[ "$OSTYPE" == cygwin ]]; then
     rtc_shared_lib_out=hiprtc-builtins64.dll
@@ -124,7 +125,7 @@ EOF
 
   set -x
 
-  $LLVM_DIR/bin/clang -O3 --rocm-path=$HIP_INC_DIR/.. -std=c++17 -nogpulib -isystem $HIP_INC_DIR -isystem $HIP_BUILD_INC_DIR --cuda-device-only -x hip $tmp/hip_pch.h -E >$tmp/pch.cui &&
+  $LLVM_DIR/bin/clang -O3 --rocm-path=$HIP_INC_DIR/.. -std=c++17 -nogpulib -isystem $HIP_INC_DIR -isystem $HIP_BUILD_INC_DIR -isystem $HIP_AMD_INC_DIR --cuda-device-only -x hip $tmp/hip_pch.h -E >$tmp/pch.cui &&
 
   cat $tmp/hip_macros.h >> $tmp/pch.cui &&
 
@@ -174,7 +175,7 @@ __hipRTC_header_size:
 EOF
 
   set -x
-  $LLVM_DIR/bin/clang -O3 --rocm-path=$HIP_INC_DIR/.. -std=c++14 -nogpulib --hip-version=4.4 -isystem $HIP_INC_DIR -isystem $HIP_BUILD_INC_DIR --cuda-device-only -D__HIPCC_RTC__ -x hip $tmp/hipRTC_header.h -E -o $tmp/hiprtc &&
+  $LLVM_DIR/bin/clang -O3 --rocm-path=$HIP_INC_DIR/.. -std=c++14 -nogpulib --hip-version=4.4 -isystem $HIP_INC_DIR -isystem $HIP_BUILD_INC_DIR -isystem $HIP_AMD_INC_DIR --cuda-device-only -D__HIPCC_RTC__ -x hip $tmp/hipRTC_header.h -E -o $tmp/hiprtc &&
   cat $macroFile >> $tmp/hiprtc &&
   $LLVM_DIR/bin/llvm-mc -o $tmp/hiprtc_header.o $tmp/hipRTC_header.mcin --filetype=obj &&
   $LLVM_DIR/bin/clang $tmp/hiprtc_header.o -o $rtc_shared_lib_out -shared &&

projects/hip/include/hip/amd_detail/cuda/cuda.h

@@ -1 +0,0 @@
-

projects/hip/include/hip/amd_detail/cuda/math_functions.h

@@ -1 +0,0 @@
-

projects/hip/include/hip/channel_descriptor.h

@@ -29,9 +29,9 @@ THE SOFTWARE.
 
 
 #if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
-#include <hip/amd_detail/channel_descriptor.h>
+#include <hip/amd_detail/amd_channel_descriptor.h>
 #elif !(defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && (defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
-#include <hip/nvidia_detail/channel_descriptor.h>
+#include <hip/nvidia_detail/nvidia_channel_descriptor.h>
 #else
 #error("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
 #endif

projects/hip/include/hip/device_functions.h

@@ -26,7 +26,7 @@ THE SOFTWARE.
 #include <hip/hip_common.h>
 
 #if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
-#include <hip/amd_detail/device_functions.h>
+#include <hip/amd_detail/amd_device_functions.h>
 #elif !(defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && (defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
 #include <device_functions.h>
 #else

projects/hip/include/hip/hip_complex.h

@@ -26,9 +26,9 @@ THE SOFTWARE.
 #include <hip/hip_common.h>
 
 #if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
-#include <hip/amd_detail/hip_complex.h>
+#include <hip/amd_detail/amd_hip_complex.h>
 #elif !(defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && (defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
-#include <hip/nvidia_detail/hip_complex.h>
+#include <hip/nvidia_detail/nvidia_hip_complex.h>
 #else
 #error("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
 #endif

projects/hip/include/hip/hip_cooperative_groups.h

@@ -35,10 +35,10 @@ THE SOFTWARE.
 
 #if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
 #if __cplusplus && defined(__clang__) && defined(__HIP__)
-#include <hip/amd_detail/hip_cooperative_groups.h>
+#include <hip/amd_detail/amd_hip_cooperative_groups.h>
 #endif
 #elif !(defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && (defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
-#include <hip/nvidia_detail/hip_cooperative_groups.h>
+#include <hip/nvidia_detail/nvidia_hip_cooperative_groups.h>
 #else
 #error("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
 #endif

projects/hip/include/hip/hip_fp16.h

@@ -26,7 +26,7 @@ THE SOFTWARE.
 #include <hip/hip_common.h>
 
 #if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
-#include <hip/amd_detail/hip_fp16.h>
+#include <hip/amd_detail/amd_hip_fp16.h>
 #elif !(defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && (defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
 #include "cuda_fp16.h"
 #else

projects/hip/include/hip/hip_runtime.h

@@ -59,9 +59,9 @@ THE SOFTWARE.
 #include <hip/hip_common.h>
 
 #if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
-#include <hip/amd_detail/hip_runtime.h>
+#include <hip/amd_detail/amd_hip_runtime.h>
 #elif !(defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && (defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
-#include <hip/nvidia_detail/hip_runtime.h>
+#include <hip/nvidia_detail/nvidia_hip_runtime.h>
 #else
 #error("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
 #endif

projects/hip/include/hip/hip_texture_types.h

@@ -25,7 +25,7 @@ THE SOFTWARE.
 #define HIP_INCLUDE_HIP_HIP_TEXTURE_TYPES_H
 
 #if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
-#include <hip/amd_detail/hip_texture_types.h>
+#include <hip/amd_detail/amd_hip_texture_types.h>
 #elif !(defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && (defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
 #include <hip/nvidia_detail/hip_texture_types.h>
 #else

projects/hip/include/hip/hip_vector_types.h

@@ -30,7 +30,7 @@ THE SOFTWARE.
 
 #if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
 #if __cplusplus
-#include <hip/amd_detail/hip_vector_types.h>
+#include <hip/amd_detail/amd_hip_vector_types.h>
 #endif
 #elif !(defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && (defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
 #include <vector_types.h>

projects/hip/include/hip/hiprtc.h

@@ -24,7 +24,7 @@ THE SOFTWARE.
 #include <hip/hip_common.h>
 
 #if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
-    #include <hip/amd_detail/hiprtc.h>
+    #include <hip/amd_detail/amd_hiprtc.h>
 #elif !(defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && (defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
     #include <hip/nvidia_detail/hiprtc.h>
 #else

projects/hip/include/hip/library_types.h

@@ -26,7 +26,7 @@ THE SOFTWARE.
 #include <hip/hip_common.h>
 
 #if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
-#include <hip/amd_detail/library_types.h>
+#include <hip/amd_detail/amd_library_types.h>
 #elif !(defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && (defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
 #include "library_types.h"
 #else

projects/hip/include/hip/math_functions.h

@@ -30,7 +30,7 @@ THE SOFTWARE.
 #include <hip/hip_common.h>
 
 #if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
-#include <hip/amd_detail/math_functions.h>
+#include <hip/amd_detail/amd_math_functions.h>
 #elif !(defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && (defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
 //#include <hip/nvidia_detail/math_functions.h>
 #else

projects/hip/include/hip/texture_types.h

@@ -26,7 +26,7 @@ THE SOFTWARE.
 #include <hip/hip_common.h>
 
 #if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
-#include <hip/amd_detail/texture_types.h>
+#include <hip/amd_detail/amd_texture_types.h>
 #elif !(defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && (defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
 #include "texture_types.h"
 #else

projects/hip/packaging/hip-rocclr.txt

@@ -31,7 +31,7 @@ endif()
 
 install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib)
 install(FILES @PROJECT_BINARY_DIR@/hip-config.cmake @PROJECT_BINARY_DIR@/hip-config-version.cmake DESTINATION lib/cmake/hip)
-install(FILES @PROJECT_BINARY_DIR@/rocclr/hip-lang-config.cmake @PROJECT_BINARY_DIR@/rocclr/hip-lang-config-version.cmake DESTINATION lib/cmake/hip-lang)
+install(FILES @PROJECT_BINARY_DIR@/src/hipamd/hip-lang-config.cmake @PROJECT_BINARY_DIR@/src/hipamd/hip-lang-config-version.cmake DESTINATION lib/cmake/hip-lang)
 
 if(@__HIP_ENABLE_RTC@)
   install(FILES @PROJECT_BINARY_DIR@/lib/libhiprtc-builtins.so.@HIP_LIB_VERSION_MAJOR@.@HIP_LIB_VERSION_MINOR@ DESTINATION lib)

projects/hip/src/hipamd/CMakeLists.txt

@@ -82,51 +82,52 @@ if(NOT WIN32)
 endif()
 
 target_sources(amdhip64 PRIVATE
-  cl_gl.cpp
-  cl_lqdflash_amd.cpp
-  fixme.cpp
-  hip_activity.cpp
-  hip_code_object.cpp
-  hip_context.cpp
-  hip_device_runtime.cpp
-  hip_device.cpp
-  hip_error.cpp
-  hip_event.cpp
-  hip_fatbin.cpp
-  hip_global.cpp
-  hip_graph_internal.cpp
-  hip_graph.cpp
-  hip_hmm.cpp
-  hip_intercept.cpp
-  hip_memory.cpp
-  hip_module.cpp
-  hip_peer.cpp
-  hip_platform.cpp
-  hip_profile.cpp
-  hip_rtc.cpp
-  hip_stream_ops.cpp
-  hip_stream.cpp
-  hip_surface.cpp
-  hip_texture.cpp)
+  src/cl_gl.cpp
+  src/cl_lqdflash_amd.cpp
+  src/fixme.cpp
+  src/hip_activity.cpp
+  src/hip_code_object.cpp
+  src/hip_context.cpp
+  src/hip_device_runtime.cpp
+  src/hip_device.cpp
+  src/hip_error.cpp
+  src/hip_event.cpp
+  src/hip_fatbin.cpp
+  src/hip_global.cpp
+  src/hip_graph_internal.cpp
+  src/hip_graph.cpp
+  src/hip_hmm.cpp
+  src/hip_intercept.cpp
+  src/hip_memory.cpp
+  src/hip_module.cpp
+  src/hip_peer.cpp
+  src/hip_platform.cpp
+  src/hip_profile.cpp
+  src/hip_rtc.cpp
+  src/hip_stream_ops.cpp
+  src/hip_stream.cpp
+  src/hip_surface.cpp
+  src/hip_texture.cpp)
 
 if(WIN32)
   target_sources(amdhip64 PRIVATE
-    cl_d3d9.cpp
-    cl_d3d10.cpp
-    cl_d3d11.cpp)
+    src/cl_d3d9.cpp
+    src/cl_d3d10.cpp
+    src/cl_d3d11.cpp)
 endif()
 
 if(BUILD_SHARED_LIBS)
   if(WIN32)
     target_sources(amdhip64 PRIVATE amdhip.def)
   else()
-    target_link_libraries(amdhip64 PRIVATE "-Wl,--version-script=${CMAKE_CURRENT_LIST_DIR}/hip_hcc.map.in")
-    set_target_properties(amdhip64 PROPERTIES LINK_DEPENDS "${CMAKE_CURRENT_LIST_DIR}/hip_hcc.map.in")
+    target_link_libraries(amdhip64 PRIVATE "-Wl,--version-script=${CMAKE_CURRENT_LIST_DIR}/src/hip_hcc.map.in")
+    set_target_properties(amdhip64 PROPERTIES LINK_DEPENDS "${CMAKE_CURRENT_LIST_DIR}/src/hip_hcc.map.in")
   endif()
 endif()
 
 target_include_directories(amdhip64
   PRIVATE
+    ${PROJECT_SOURCE_DIR}/src/hipamd/include
     ${PROJECT_SOURCE_DIR}/include
     ${PROJECT_BINARY_DIR}/include)
 
@@ -156,7 +157,8 @@ if(__HIP_ENABLE_PCH)
   # find_package(LLVM) returns the lib/cmake/llvm location. We require the root.
   set(HIP_LLVM_ROOT "${LLVM_DIR}/../../..")
 
-  execute_process(COMMAND sh -c "${CMAKE_CURRENT_SOURCE_DIR}/../bin/hip_embed_pch.sh ${PROJECT_BINARY_DIR}/include ${PROJECT_SOURCE_DIR}/include ${HIP_LLVM_ROOT}" COMMAND_ECHO STDERR RESULT_VARIABLE EMBED_PCH_RC)
+  #  execute_process(COMMAND sh -c "${CMAKE_CURRENT_SOURCE_DIR}/../bin/hip_embed_pch.sh ${PROJECT_BINARY_DIR}/include ${PROJECT_SOURCE_DIR}/include ${HIP_LLVM_ROOT}" COMMAND_ECHO STDERR RESULT_VARIABLE EMBED_PCH_RC)
+  execute_process(COMMAND sh -c "${CMAKE_CURRENT_SOURCE_DIR}/../../bin/hip_embed_pch.sh ${PROJECT_BINARY_DIR}/include ${PROJECT_SOURCE_DIR}/include ${PROJECT_SOURCE_DIR}/src/hipamd/include ${HIP_LLVM_ROOT}" COMMAND_ECHO STDERR RESULT_VARIABLE EMBED_PCH_RC)
   if (EMBED_PCH_RC AND NOT EMBED_PCH_RC EQUAL 0)
     message(FATAL_ERROR "Failed to embed PCH")
   endif()
@@ -179,7 +181,7 @@ if(__HIP_ENABLE_RTC)
     set(HIPRTC_LIB_NAME "libhiprtc-builtins.so.${HIP_LIB_VERSION_MAJOR}.${HIP_LIB_VERSION_MINOR}")
   endif()
   execute_process(
-    COMMAND sh -c "mkdir -p ${PROJECT_BINARY_DIR}/lib; ${CMAKE_CURRENT_SOURCE_DIR}/../bin/hip_embed_pch.sh ${PROJECT_BINARY_DIR}/include ${PROJECT_SOURCE_DIR}/include ${HIP_LLVM_ROOT} -r ${PROJECT_BINARY_DIR}/lib/${HIPRTC_LIB_NAME}"
+    COMMAND sh -c "mkdir -p ${PROJECT_BINARY_DIR}/lib; ${CMAKE_CURRENT_SOURCE_DIR}/../../bin/hip_embed_pch.sh ${PROJECT_BINARY_DIR}/include ${PROJECT_SOURCE_DIR}/include ${PROJECT_SOURCE_DIR}/src/hipamd/include ${HIP_LLVM_ROOT} -r ${PROJECT_BINARY_DIR}/lib/${HIPRTC_LIB_NAME}"
     COMMAND_ECHO STDERR
     RESULT_VARIABLE EMBED_RTC_RC
   )
@@ -194,9 +196,9 @@ endif()
 #############################
 # Generate profiling API macros/structures header
 set(PROF_API_STR "${PROJECT_BINARY_DIR}/include/hip/amd_detail/hip_prof_str.h")
-set(PROF_API_HDR "${PROJECT_SOURCE_DIR}/include/hip/amd_detail/hip_runtime_api.h")
-set(PROF_API_SRC "${CMAKE_CURRENT_SOURCE_DIR}")
-set(PROF_API_GEN "${CMAKE_CURRENT_SOURCE_DIR}/hip_prof_gen.py")
+set(PROF_API_HDR "${PROJECT_SOURCE_DIR}/include/hip/hip_runtime_api.h")
+set(PROF_API_SRC "${CMAKE_CURRENT_SOURCE_DIR}/src")
+set(PROF_API_GEN "${CMAKE_CURRENT_SOURCE_DIR}/src/hip_prof_gen.py")
 set(PROF_API_LOG "${PROJECT_BINARY_DIR}/hip_prof_gen.log.txt")
 
 find_package(PythonInterp REQUIRED)

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

@@ -24,8 +24,8 @@ THE SOFTWARE.
 #define HIP_INCLUDE_HIP_AMD_DETAIL_CHANNEL_DESCRIPTOR_H
 
 #include <hip/hip_common.h>
-#include <hip/amd_detail/driver_types.h>
-#include <hip/amd_detail/hip_vector_types.h>
+#include <hip/amd_detail/amd_driver_types.h>
+#include <hip/amd_detail/amd_hip_vector_types.h>
 
 #ifdef __cplusplus
 

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

@@ -1,6 +1,6 @@
 
 
-#include "device_functions.h"
+#include "amd_device_functions.h"
 
 #if __has_builtin(__hip_atomic_compare_exchange_strong)
 

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

@@ -23,7 +23,7 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
 #define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
 
-#include "hip/amd_detail/hip_vector_types.h"
+#include "hip/amd_detail/amd_hip_vector_types.h"
 
 #if defined(__HIPCC_RTC__)
 #define __HOST_DEVICE__ __device__

projects/hip/src/hipamd/include/hip/amd_detail/amd_hip_fp16.h

@@ -24,7 +24,8 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_FP16_H
 #define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_FP16_H
 
-#include <hip/amd_detail/hip_common.h>
+#include <hip/amd_detail/amd_hip_common.h>
+
 #include "hip/amd_detail/host_defines.h"
 #if defined(__HIPCC_RTC__)
   #define __HOST_DEVICE__ __device__
@@ -64,7 +65,7 @@ THE SOFTWARE.
 
     #if defined(__cplusplus)
         #include "hip_fp16_math_fwd.h"
-        #include "hip_vector_types.h"
+        #include "amd_hip_vector_types.h"
         #include "host_defines.h"
 
         namespace std

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

@@ -29,7 +29,7 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_RUNTIME_H
 #define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_RUNTIME_H
 
-#include <hip/amd_detail/hip_common.h>
+#include <hip/amd_detail/amd_hip_common.h>
 
 //---
 // Top part of file can be compiled with any compiler
@@ -67,10 +67,10 @@ extern int HIP_TRACE_API;
 #ifdef __cplusplus
 #include <hip/amd_detail/hip_ldg.h>
 #endif
-#include <hip/amd_detail/hip_atomic.h>
+#include <hip/amd_detail/amd_hip_atomic.h>
 #include <hip/amd_detail/host_defines.h>
-#include <hip/amd_detail/device_functions.h>
-#include <hip/amd_detail/surface_functions.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>
 
@@ -359,7 +359,7 @@ extern "C" __device__ __attribute__((const)) size_t __ockl_get_num_groups(uint);
 #define hipGridDim_y (__ockl_get_num_groups(1))
 #define hipGridDim_z (__ockl_get_num_groups(2))
 
-#include <hip/amd_detail/math_functions.h>
+#include <hip/amd_detail/amd_math_functions.h>
 
 #if __HIP_HCC_COMPAT_MODE__
 // Define HCC work item functions in terms of HIP builtin variables.

projects/hip/src/hipamd/include/hip/amd_detail/amd_hip_surface_types.h

@@ -28,7 +28,7 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_SURFACE_TYPES_H
 #define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_SURFACE_TYPES_H
 
-#include <hip/amd_detail/driver_types.h>
+#include <hip/amd_detail/amd_driver_types.h>
 
 /**
  * An opaque value that represents a hip surface object

projects/hip/src/hipamd/include/hip/amd_detail/amd_hip_texture_types.h

@@ -35,9 +35,9 @@ THE SOFTWARE.
  *******************************************************************************/
 #if !defined(__HIPCC_RTC__)
 #include <limits.h>
-#include <hip/amd_detail/channel_descriptor.h>
+#include <hip/amd_detail/amd_channel_descriptor.h>
 #endif // !defined(__HIPCC_RTC__)
-#include <hip/amd_detail/texture_types.h>
+#include <hip/amd_detail/amd_texture_types.h>
 
 #if __cplusplus
 

projects/hip/src/hipamd/include/hip/amd_detail/amd_math_functions.h

@@ -23,7 +23,7 @@ THE SOFTWARE.
 #pragma once
 
 #include "hip_fp16_math_fwd.h"
-#include "hip_vector_types.h"
+#include "amd_hip_vector_types.h"
 #include "math_fwd.h"
 
 #include <hip/amd_detail/host_defines.h>
@@ -1499,4 +1499,4 @@ inline _Float16 pow(_Float16 base, int iexp) {
 // For backward compatibility.
 // There are HIP applications e.g. TensorFlow, expecting __HIP_ARCH_* macros
 // defined after including math_functions.h.
-#include <hip/amd_detail/hip_runtime.h>
+#include <hip/amd_detail/amd_hip_runtime.h>

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

@@ -23,7 +23,7 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_SURFACE_FUNCTIONS_H
 #define HIP_INCLUDE_HIP_AMD_DETAIL_SURFACE_FUNCTIONS_H
 
-#include <hip/amd_detail/hip_surface_types.h>
+#include <hip/amd_detail/amd_hip_surface_types.h>
 
 #define __SURFACE_FUNCTIONS_DECL__ static inline __device__
 template <class T>

projects/hip/src/hipamd/include/hip/amd_detail/amd_texture_types.h

@@ -24,7 +24,7 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_TEXTURE_TYPES_H
 #define HIP_INCLUDE_HIP_AMD_DETAIL_TEXTURE_TYPES_H
 
-#include <hip/amd_detail/driver_types.h>
+#include <hip/amd_detail/amd_driver_types.h>
 
 #define hipTextureType1D 0x01
 #define hipTextureType2D 0x02

projects/hip/src/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h

@@ -32,8 +32,7 @@ THE SOFTWARE.
 #define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_HELPER_H
 
 #if __cplusplus
-#include <hip/amd_detail/hip_runtime_api.h>
-#include <hip/amd_detail/device_functions.h>
+#include <hip/amd_detail/amd_device_functions.h>
 
 #if !defined(__align__)
 #define __align__(x) __attribute__((aligned(x)))

projects/hip/src/hipamd/include/hip/amd_detail/hip_ldg.h

@@ -24,7 +24,7 @@ THE SOFTWARE.
 #define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_LDG_H
 
 #if __HIP_CLANG_ONLY__
-#include "hip_vector_types.h"
+#include "amd_hip_vector_types.h"
 #include "host_defines.h"
 
 __device__ inline static char __ldg(const char* ptr) { return *ptr; }

projects/hip/src/hipamd/include/hip/amd_detail/texture_functions.h

@@ -22,8 +22,8 @@ THE SOFTWARE.
 
 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_TEXTURE_FUNCTIONS_H
 #define HIP_INCLUDE_HIP_AMD_DETAIL_TEXTURE_FUNCTIONS_H
-#include <hip/amd_detail/hip_vector_types.h>
-#include <hip/amd_detail/hip_texture_types.h>
+#include <hip/amd_detail/amd_hip_vector_types.h>
+#include <hip/amd_detail/amd_hip_texture_types.h>
 
 #pragma push_macro("TYPEDEF_VECTOR_VALUE_TYPE")
 #define TYPEDEF_VECTOR_VALUE_TYPE(SCALAR_TYPE) \

projects/hip/src/hipamd/include/hip/nvcc_detail

@@ -0,0 +1 @@
+nvidia_detail

projects/hip/src/hipamd/src/hip_code_object.cpp

@@ -24,7 +24,7 @@ THE SOFTWARE.
 
 #include <cstring>
 
-#include <hip/amd_detail/driver_types.h>
+#include <hip/amd_detail/amd_driver_types.h>
 #include "hip/hip_runtime_api.h"
 #include "hip/hip_runtime.h"
 #include "hip_internal.hpp"

projects/hip/src/hipamd/src/hip_conversions.hpp

@@ -22,8 +22,8 @@ THE SOFTWARE.
 
 #pragma once
 
-#include <hip/amd_detail/driver_types.h>
-#include <hip/amd_detail/texture_types.h>
+#include <hip/amd_detail/amd_driver_types.h>
+#include <hip/amd_detail/amd_texture_types.h>
 
 namespace hip
 {