AI/rocm-systems
2
0
Fork 0
Koodi Ongelmat Pull-pyynnöt Actions Paketit Projektit Julkaisut Wiki Toiminta

SWDEV-249870 - Short-Term solution for Pre-Compiled Headers for Online Compilation

Selaa lähdekoodia 
Change-Id: Iec989787e546ff2eb68c4b146dc540655698b569
This commit is contained in:
agodavar
2020-09-18 03:33:41 -04:00
committed by Anusha Godavarthy Surya
vanhempi eedf4c68dc
commit 3c9bbc335a
7 muutettua tiedostoa jossa 36 lisäystä ja 54 poistoa
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/CMakeLists.txt
+8 -2
Näytä tiedosto
@@ -8,10 +8,15 @@ set(BUILD_SHARED_LIBS ON CACHE BOOL "Build shared library (.so) or static lib (
set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
if(NOT ${BUILD_SHARED_LIBS} AND NOT DEFINED ENABLE_HIP_PCH)
set(ENABLE_HIP_PCH ON CACHE BOOL "enable/disable pre-compiled hip headers")
if(NOT DEFINED __HIP_ENABLE_PCH)
set(__HIP_ENABLE_PCH ON CACHE BOOL "enable/disable pre-compiled hip headers")
endif()
if(${__HIP_ENABLE_PCH})
set(_pchStatus 1)
else()
set(_pchStatus 0)
endif()
#############################
# Options
#############################
@@ -455,6 +460,7 @@ set(_versionInfoHeader
#define HIP_VERSION_MINOR ${HIP_VERSION_MINOR}
#define HIP_VERSION_PATCH ${HIP_VERSION_GITDATE}
#define HIP_VERSION (HIP_VERSION_MAJOR * 100 + HIP_VERSION_MINOR)\n
#define __HIP_HAS_GET_PCH ${_pchStatus}\n
#endif\n
")
file(WRITE "${PROJECT_BINARY_DIR}/include/hip/hip_version.h" ${_versionInfoHeader})
hipamd/bin/hip_embed_pch.sh
+4 -5
Näytä tiedosto
@@ -1,8 +1,7 @@
#!/bin/bash
#set -x
ROCM_PATH=${ROCM_PATH:-/opt/rocm}
LLVM_DIR="$1/../../../"
tmp=/tmp/hip_pch.$$
mkdir -p $tmp
@@ -47,12 +46,12 @@ __hip_pch_size:
.long __hip_pch_size - __hip_pch
EOF
$ROCM_PATH/llvm/bin/clang -O3 -c -std=c++17 -isystem /opt/rocm/llvm/lib/clang/11.0.0/include/.. -isystem /opt/rocm/include -nogpulib --cuda-device-only -x hip $tmp/hip_pch.h -E >$tmp/pch.cui
$LLVM_DIR/bin/clang -O3 -c -std=c++17 -isystem $LLVM_DIR/lib/clang/11.0.0/include/.. -isystem /opt/rocm/include -nogpulib --cuda-device-only -x hip $tmp/hip_pch.h -E >$tmp/pch.cui
cat $tmp/hip_macros.h >> $tmp/pch.cui
$ROCM_PATH/llvm/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
$ROCM_PATH/llvm/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
hipamd/bin/hip_gen_pch.sh
-36
Näytä tiedosto
@@ -1,36 +0,0 @@
#!/bin/bash
#set -x
cat >/tmp/hip_macros.h <<EOF
#define __device__ __attribute__((device))
#define __host__ __attribute__((host))
#define __global__ __attribute__((global))
#define __constant__ __attribute__((constant))
#define __shared__ __attribute__((shared))
#define launch_bounds_impl0(requiredMaxThreadsPerBlock) \
__attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock)))
#define launch_bounds_impl1(requiredMaxThreadsPerBlock, minBlocksPerMultiprocessor) \
__attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock), \
amdgpu_waves_per_eu(minBlocksPerMultiprocessor)))
#define select_impl_(_1, _2, impl_, ...) impl_
#define __launch_bounds__(...) \
select_impl_(__VA_ARGS__, launch_bounds_impl1, launch_bounds_impl0)(__VA_ARGS__)
// Macro to replace extern __shared__ declarations
// to local variable definitions
#define HIP_DYNAMIC_SHARED(type, var) \
type* var = (type*)__amdgcn_get_dynamicgroupbaseptr();
EOF
cat >/tmp/hip_pch.h <<EOF
#include "hip/hip_runtime.h"
#include "hip/hip_fp16.h"
EOF
/opt/rocm/llvm/bin/clang -O3 -c -std=c++17 -isystem /opt/rocm/llvm/lib/clang/11.0.0/include/.. -isystem /opt/rocm/include -nogpulib --cuda-device-only -x hip /tmp/hip_pch.h -E >/tmp/pch.cui
cat /tmp/hip_macros.h >> /tmp/pch.cui
/opt/rocm/llvm/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
hipamd/include/hip/hcc_detail/hip_runtime_api.h
+8 -5
Näytä tiedosto
@@ -345,13 +345,16 @@ typedef struct hipLaunchParams_t {
hipStream_t stream; ///< Stream identifier
} hipLaunchParams;
// Pre-Compiled header for online compilation
#ifdef ENABLE_HIP_PCH
extern const char* __hip_pch;
extern unsigned __hip_pch_size;
void __hipGetPCH(const char** pch, unsigned int*size);
#if __HIP_HAS_GET_PCH
/**
* Internal use only. This API may change in the future
* Pre-Compiled header for online compilation
*
*/
void __hipGetPCH(const char** pch, unsigned int*size);
#endif
// Doxygen end group GlobalDefs
/** @} */
hipamd/include/hip/hip_runtime_api.h
+1
Näytä tiedosto
@@ -32,6 +32,7 @@ THE SOFTWARE.
#include <string.h> // for getDeviceProp
#include <hip/hip_version.h>
#include <hip/hip_common.h>
enum {
hipamd/rocclr/CMakeLists.txt
+12 -5
Näytä tiedosto
@@ -96,6 +96,14 @@ find_package(amd_comgr REQUIRED CONFIG
message(STATUS "Code Object Manager found at ${amd_comgr_DIR}.")
find_package(LLVM REQUIRED CONFIG
PATHS
/opt/rocm/llvm
PATH_SUFFIXES
lib/cmake/llvm)
message(STATUS "llvm found at ${LLVM_DIR}.")
add_library(hip64 OBJECT
hip_context.cpp
hip_code_object.cpp
@@ -148,10 +156,9 @@ endif()
# Short-Term solution for pre-compiled headers for online compilation
# Enable pre compiled header
if(${ENABLE_HIP_PCH})
execute_process(COMMAND sh -c "${CMAKE_CURRENT_SOURCE_DIR}/../bin/hip_gen_pch.sh")
execute_process(COMMAND sh -c "${CMAKE_CURRENT_SOURCE_DIR}/../bin/hip_embed_pch.sh")
add_definitions(-DENABLE_HIP_PCH)
if(${__HIP_ENABLE_PCH})
execute_process(COMMAND sh -c "${CMAKE_CURRENT_SOURCE_DIR}/../bin/hip_embed_pch.sh ${LLVM_DIR}")
add_definitions(-D__HIP_ENABLE_PCH)
endif()
# Enable profiling API
@@ -216,7 +223,7 @@ add_library(device INTERFACE)
target_link_libraries(device INTERFACE host)
# Short-Term solution for pre-compiled headers for online compilation
if(${ENABLE_HIP_PCH})
if(${__HIP_ENABLE_PCH})
target_link_libraries(amdhip64 PRIVATE ${CMAKE_BINARY_DIR}/hip_pch.o)
endif()
hipamd/rocclr/hip_global.cpp
+3 -1
Näytä tiedosto
@@ -5,7 +5,9 @@
#include "hip_code_object.hpp"
#include "platform/program.hpp"
#ifdef ENABLE_HIP_PCH
#ifdef __HIP_ENABLE_PCH
extern const char __hip_pch[];
extern unsigned __hip_pch_size;
void __hipGetPCH(const char** pch, unsigned int *size) {
*pch = __hip_pch;
*size = __hip_pch_size;