From 2e6662b773bd6bfe282beef94f4efa46578f0aa5 Mon Sep 17 00:00:00 2001
From: Aaron En Ye Shi <enye.shi@gmail.com>
Date: Tue, 1 Jun 2021 17:56:42 +0000
Subject: [PATCH] SWDEV-283266 - Fix hipRTC script bugs

Remove extra -isystem causing --cuda-device-only flag
to be dropped. Fix missing __hipRTC_header_size symbol.
Add macros which were consumed during pre-processor
generation. Add INT_MAX and CHAR_BIT macros used in
HIP headers.

Change-Id: Id5143e3c8a2b1e7c78658ba84d5ab3b55ac1fa72


[ROCm/clr commit: 7c8e210c26665af3a8ec82ab6dc93c711540e7a5]
---
 projects/clr/hipamd/bin/hip_embed_pch.sh | 31 +++++++++++++++++++-----
 1 file changed, 25 insertions(+), 6 deletions(-)

diff --git a/projects/clr/hipamd/bin/hip_embed_pch.sh b/projects/clr/hipamd/bin/hip_embed_pch.sh
index cc4f87a49a..0ea905a268 100755
--- a/projects/clr/hipamd/bin/hip_embed_pch.sh
+++ b/projects/clr/hipamd/bin/hip_embed_pch.sh
@@ -77,11 +77,9 @@ else
   tmpdir=/tmp
 fi
 
-generate_pch() {
-  tmp=$tmpdir/hip_pch.$$
-  mkdir -p $tmp
-
-cat >$tmp/hip_macros.h <<EOF
+# Expected first argument $1 to be output file name.
+create_hip_macro_file() {
+cat >$1 <<EOF
 #define __device__ __attribute__((device))
 #define __host__ __attribute__((host))
 #define __global__ __attribute__((global))
@@ -98,6 +96,13 @@ cat >$tmp/hip_macros.h <<EOF
     select_impl_(__VA_ARGS__, launch_bounds_impl1, launch_bounds_impl0)(__VA_ARGS__)
 
 EOF
+}
+
+generate_pch() {
+  tmp=$tmpdir/hip_pch.$$
+  mkdir -p $tmp
+
+  create_hip_macro_file $tmp/hip_macros.h
 
 cat >$tmp/hip_pch.h <<EOF
 #include "hip/hip_runtime.h"
@@ -133,17 +138,29 @@ EOF
 generate_rtc_header() {
   tmp=$tmpdir/hip_rtc.$$
   mkdir -p $tmp
+  local macroFile="$tmp/hip_macros.h"
   local headerFile="$tmp/hipRTC_header.h"
   local mcinFile="$tmp/hipRTC_header.mcin"
 
+  create_hip_macro_file $macroFile
+
 cat >$headerFile <<EOF
+#pragma push_macro("CHAR_BIT")
+#pragma push_macro("INT_MAX")
+#define CHAR_BIT __CHAR_BIT__
+#define INT_MAX __INTMAX_MAX__
+
 #include "hip/hip_runtime.h"
 #include "hip/hip_fp16.h"
+
+#pragma pop_macro("CHAR_BIT")
+#pragma pop_macro("INT_MAX")
 EOF
 
   echo "// Automatically generated script for HIP RTC." > $mcinFile
   if [[ $isWindows -eq 0 ]]; then
     echo "  .type __hipRTC_header,@object" >> $mcinFile
+    echo "  .type __hipRTC_header_size,@object" >> $mcinFile
   fi
 cat >>$mcinFile <<EOF
   .section .hipRTC_header,"a"
@@ -157,9 +174,11 @@ __hipRTC_header_size:
 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 -isystem --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 --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 &&
+  $LLVM_DIR/bin/clang -O3 --rocm-path=$HIP_INC_DIR/.. -std=c++14 -nogpulib -nogpuinc -emit-llvm -c -o $tmp/tmp.bc --cuda-device-only -D__HIPCC_RTC__ --offload-arch=gfx906 -x hip-cpp-output $tmp/hiprtc &&
   rm -rf $tmp
 }