From 4dedd2436d6662e8d7611bb4abcd965452131f52 Mon Sep 17 00:00:00 2001
From: Aaron En Ye Shi <enye.shi@gmail.com>
Date: Thu, 6 May 2021 15:44:44 +0000
Subject: [PATCH] SWDEV-283266 - Add hiprtc option to generate preprocessor
 expansion

Add -r/--generate_rtc option to hip_embed_pch.sh to generate the
preprocessor expansion output of the HIP headers for hipRTC mode.
Generates an object hiprtc_header.o and shared library libhiprtc.so
which can be used for hipRTC online compilations. Enable
__HIP_ENABLE_RTC by default in the CMake files.

Change-Id: Ief51b695c9b31941d929e06f16872f95eea3619b
---
 CMakeLists.txt           |   1 +
 bin/hip_embed_pch.sh     | 114 ++++++++++++++++++++++++++++++++++-----
 packaging/hip-rocclr.txt |   4 ++
 rocclr/CMakeLists.txt    |  21 +++++++-
 4 files changed, 126 insertions(+), 14 deletions(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index f97fb8395f..76fa45d8ce 100755
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -42,6 +42,7 @@ set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
 #############################
 option(BUILD_HIPIFY_CLANG "Enable building the CUDA->HIP converter" OFF)
 option(__HIP_ENABLE_PCH "Enable/Disable pre-compiled hip headers" ON)
+option(__HIP_ENABLE_RTC "Enable/Disable pre-processed hiprtc shared lib" ON)
 
 if(__HIP_ENABLE_PCH)
   set(_pchStatus 1)
diff --git a/bin/hip_embed_pch.sh b/bin/hip_embed_pch.sh
index c77d084783..a4cb64c972 100755
--- a/bin/hip_embed_pch.sh
+++ b/bin/hip_embed_pch.sh
@@ -19,13 +19,66 @@
 # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 # THE SOFTWARE.
 
-#set -x
-HIP_BUILD_INC_DIR=$1
-HIP_INC_DIR=$2
-LLVM_DIR="$3/../../../"
+printUsage() {
+  echo
+  echo "Usage: $(basename "$0") HIP_BUILD_INC_DIR HIP_INC_DIR LLVM_DIR HSA_DIR [option] [RTC_LIB_OUTPUT]"
+  echo
+  echo "Options:"
+  echo "  -p,  --generate_pch  Generate pre-compiled header (default)"
+  echo "  -r,  --generate_rtc  Generate preprocessor expansion (hiprtc_header.o)"
+  echo "  -h,  --help          Prints this help"
+  echo
+  echo
+  return 0
+}
+
+if [ "$1" == "" ]; then
+  printUsage
+  exit 0
+fi
+
+HIP_BUILD_INC_DIR="$1"
+HIP_INC_DIR="$2"
+LLVM_DIR="$3"
 HSA_DIR="$4"
-tmp=/tmp/hip_pch.$$
-mkdir -p $tmp
+# By default, generate pch
+TARGET="generatepch"
+
+while [ "$5" != "" ];
+do
+  case "$5" in
+    -h | --help )
+        printUsage ; exit 0 ;;
+    -p | --generate_pch )
+        TARGET="generatepch" ; break ;;
+    -r | --generate_rtc )
+        TARGET="generatertc" ; break ;;
+    *)
+        echo " UNEXPECTED ERROR Parm : [$5] ">&2 ; exit 20 ;;
+  esac
+  shift 1
+done
+
+# Allow hiprtc lib name to be set by argument 6
+if [[ "$6" != "" ]]; then
+  rtc_shared_lib_out="$6"
+else
+  if [[ "$OSTYPE" == cygwin ]]; then
+    rtc_shared_lib_out=hiprtc-builtins64.dll
+  else
+    rtc_shared_lib_out=libhiprtc-builtins.so
+  fi
+fi
+
+if [[ "$OSTYPE" == cygwin || "$OSTYPE" == msys ]]; then
+  tmpdir=.
+else
+  tmpdir=/tmp
+fi
+
+generate_pch() {
+  tmp=$tmpdir/hip_pch.$$
+  mkdir -p $tmp
 
 cat >$tmp/hip_macros.h <<EOF
 #define __device__ __attribute__((device))
@@ -51,7 +104,7 @@ cat >$tmp/hip_pch.h <<EOF
 EOF
 
 cat >$tmp/hip_pch.mcin <<EOF
- .type __hip_pch,@object
+  .type __hip_pch,@object
   .section .hip_pch,"aMS",@progbits,1
   .data
   .globl __hip_pch
@@ -63,14 +116,49 @@ __hip_pch_size:
   .long __hip_pch_size - __hip_pch
 EOF
 
-set -x
+  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 $HSA_DIR/include --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 $HSA_DIR/include --cuda-device-only -x hip $tmp/hip_pch.h -E >$tmp/pch.cui &&
 
-cat $tmp/hip_macros.h >> $tmp/pch.cui &&
+  cat $tmp/hip_macros.h >> $tmp/pch.cui &&
 
-$LLVM_DIR/bin/clang -cc1 -O3 -emit-pch -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device -std=c++17 -fgnuc-version=4.2.1 -o $tmp/hip.pch -x hip-cpp-output - <$tmp/pch.cui &&
+  $LLVM_DIR/bin/clang -cc1 -O3 -emit-pch -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device -std=c++17 -fgnuc-version=4.2.1 -o $tmp/hip.pch -x hip-cpp-output - <$tmp/pch.cui &&
 
-$LLVM_DIR/bin/llvm-mc -o hip_pch.o $tmp/hip_pch.mcin --filetype=obj &&
+  $LLVM_DIR/bin/llvm-mc -o hip_pch.o $tmp/hip_pch.mcin --filetype=obj &&
+
+  rm -rf $tmp
+}
+
+generate_rtc_header() {
+  tmp=$tmpdir/hip_rtc.$$
+  mkdir -p $tmp
+
+cat >$tmp/hipRTC_header.h <<EOF
+#include "hip/hip_runtime.h"
+#include "hip/hip_fp16.h"
+EOF
+
+cat >$tmp/hipRTC_header.mcin <<EOF
+  .section .hipRTC_header,"a"
+  .globl __hipRTC_header
+  .globl __hipRTC_header_size
+  .p2align 3
+__hipRTC_header:
+  .incbin "$tmp/hiprtc"
+__hipRTC_header_size:
+  .long __hipRTC_header_size - __hipRTC_header
+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/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 &&
+  rm -rf $tmp
+}
+
+case $TARGET in
+    (generatertc) generate_rtc_header ;;
+    (generatepch) generate_pch ;;
+    (*) die "Invalid target $TARGET" ;;
+esac
 
-rm -rf $tmp
diff --git a/packaging/hip-rocclr.txt b/packaging/hip-rocclr.txt
index a202f9c5e0..1a8a38f5db 100644
--- a/packaging/hip-rocclr.txt
+++ b/packaging/hip-rocclr.txt
@@ -33,6 +33,10 @@ 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)
 
+if(@__HIP_ENABLE_RTC@)
+  install(FILES @PROJECT_BINARY_DIR@/lib/libhiprtc-builtins.so.@HIP_LIB_VERSION_MAJOR@.@HIP_LIB_VERSION_MINOR@ DESTINATION lib)
+endif()
+
 #############################
 # Rocclr install
 #############################
diff --git a/rocclr/CMakeLists.txt b/rocclr/CMakeLists.txt
index 63229359d0..e503b1a205 100755
--- a/rocclr/CMakeLists.txt
+++ b/rocclr/CMakeLists.txt
@@ -116,6 +116,7 @@ find_package(LLVM REQUIRED CONFIG
      lib/cmake/llvm)
 
 message(STATUS "llvm found at ${LLVM_DIR}.")
+set(LLVM_ROOT "${LLVM_DIR}/../../..")
 
 add_library(hip64 OBJECT
  hip_context.cpp
@@ -174,13 +175,31 @@ endif()
 # Short-Term solution for pre-compiled headers for online compilation
 # Enable pre compiled header
 if(__HIP_ENABLE_PCH)
-    execute_process(COMMAND sh -c "${CMAKE_CURRENT_SOURCE_DIR}/../bin/hip_embed_pch.sh ${PROJECT_BINARY_DIR}/include ${PROJECT_SOURCE_DIR}/include ${LLVM_DIR} ${HSA_PATH}" 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 ${LLVM_ROOT} ${HSA_PATH}" 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()
     add_definitions(-D__HIP_ENABLE_PCH)
 endif()
 
+# Enable preprocessed hiprtc-builtins library
+if(__HIP_ENABLE_RTC)
+  if(WIN32)
+    set(HIPRTC_LIB_NAME "hiprtc-builtins64_${HIP_LIB_VERSION_MAJOR}${HIP_LIB_VERSION_MINOR}.dll")
+  else()
+    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 ${LLVM_ROOT} ${HSA_PATH} -r ${PROJECT_BINARY_DIR}/lib/${HIPRTC_LIB_NAME}"
+    COMMAND_ECHO STDERR
+    RESULT_VARIABLE EMBED_RTC_RC
+  )
+  if (EMBED_RTC_RC AND NOT EMBED_RTC_RC EQUAL 0)
+    message(FATAL_ERROR "Failed to create hiprtc shared lib")
+  endif()
+  install(FILES ${PROJECT_BINARY_DIR}/lib/${HIPRTC_LIB_NAME} DESTINATION lib)
+endif()
+
 # Enable profiling API
 if(USE_PROF_API EQUAL 1)
   find_path(PROF_API_HEADER_DIR prof_protocol.h