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
This commit is contained in:
Aaron En Ye Shi
committed by
Aaron En Ye Shi
Aaron En Ye Shi
parent
fac2a9fbc8
commit
4dedd2436d
+101
-13
@@ -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
|
||||
|
||||
Reference in New Issue
Block a user