From 842b9d0b56ca64ea68dbbc7a0cc7980f5e568c9c Mon Sep 17 00:00:00 2001
From: Satyanvesh Dittakavi <Satyanvesh.Dittakavi@amd.com>
Date: Wed, 11 Jan 2023 15:24:19 +0000
Subject: [PATCH] SWDEV-373231 - Support hip atomics header with HIPRTC

Change-Id: Ic500dfb66ed5ef0c4b672dfd0595409d0e8caaf5
---
 hipamd/include/hip/amd_detail/amd_hip_atomic.h | 2 ++
 hipamd/src/hiprtc/CMakeLists.txt               | 3 ++-
 hipamd/src/hiprtc/cmake/HIPRTC.cmake           | 2 ++
 3 files changed, 6 insertions(+), 1 deletion(-)

diff --git a/hipamd/include/hip/amd_detail/amd_hip_atomic.h b/hipamd/include/hip/amd_detail/amd_hip_atomic.h
index 6378012beb..869f495c89 100644
--- a/hipamd/include/hip/amd_detail/amd_hip_atomic.h
+++ b/hipamd/include/hip/amd_detail/amd_hip_atomic.h
@@ -35,7 +35,9 @@ THE SOFTWARE.
 #define __HIP_MEMORY_SCOPE_SYSTEM 5
 #endif
 
+#if !defined(__HIPCC_RTC__)
 #include "amd_hip_unsafe_atomics.h"
+#endif
 
 __device__
 inline
diff --git a/hipamd/src/hiprtc/CMakeLists.txt b/hipamd/src/hiprtc/CMakeLists.txt
index e8cc47f8cc..e2c409d511 100644
--- a/hipamd/src/hiprtc/CMakeLists.txt
+++ b/hipamd/src/hiprtc/CMakeLists.txt
@@ -132,6 +132,7 @@ 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_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")
 
 # Generate required HIPRTC files.
 FILE(MAKE_DIRECTORY ${HIPRTC_GEN_DIR})
@@ -144,7 +145,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_PREPROCESSED_FILE=${HIPRTC_GEN_PREPROCESSED} -P ${HIPRTC_CMAKE}
+  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_PREPROCESSED_FILE=${HIPRTC_GEN_PREPROCESSED} -P ${HIPRTC_CMAKE}
   DEPENDS clang ${HIPRTC_GEN_HEADER})
 add_custom_command(
   OUTPUT ${HIPRTC_GEN_OBJ}
diff --git a/hipamd/src/hiprtc/cmake/HIPRTC.cmake b/hipamd/src/hiprtc/cmake/HIPRTC.cmake
index 41e02b5f48..19b59f3655 100644
--- a/hipamd/src/hiprtc/cmake/HIPRTC.cmake
+++ b/hipamd/src/hiprtc/cmake/HIPRTC.cmake
@@ -67,6 +67,8 @@ if(HIPRTC_ADD_MACROS)
   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}")
 endif()
 
 macro(generate_hiprtc_header HiprtcHeader)