Refactor for building HIP as dynamic library
Change-Id: I65a3d9d589c4fdbbdcf1611e5427224253be8260
Этот коммит содержится в:
Maneesh Gupta
+15
-8
@@ -173,33 +173,38 @@ if(HIP_PLATFORM STREQUAL "hcc")
|
||||
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${HIP_HCC_BUILD_FLAGS}")
|
||||
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${HIP_HCC_BUILD_FLAGS}")
|
||||
|
||||
set(SOURCE_FILES src/device_util.cpp
|
||||
set(SOURCE_FILES_SHARED
|
||||
src/hip_hcc.cpp
|
||||
src/hip_context.cpp
|
||||
src/hip_device.cpp
|
||||
src/hip_error.cpp
|
||||
src/hip_event.cpp
|
||||
src/hip_ldg.cpp
|
||||
src/hip_memory.cpp
|
||||
src/hip_peer.cpp
|
||||
src/hip_stream.cpp
|
||||
src/hip_fp16.cpp
|
||||
src/hip_module.cpp)
|
||||
|
||||
set(SOURCE_FILES_STATIC
|
||||
src/device_util.cpp
|
||||
src/hip_ldg.cpp
|
||||
src/hip_fp16.cpp)
|
||||
|
||||
if(${HIP_LIB_TYPE} EQUAL 0)
|
||||
add_library(hip_hcc OBJECT ${SOURCE_FILES})
|
||||
add_library(hip_hcc OBJECT ${SOURCE_FILES_SHARED} ${SOURCE_FILES_STATIC})
|
||||
elseif(${HIP_LIB_TYPE} EQUAL 1)
|
||||
add_library(hip_hcc STATIC ${SOURCE_FILES})
|
||||
add_library(hip_hcc STATIC ${SOURCE_FILES_SHARED} ${SOURCE_FILES_STATIC})
|
||||
else()
|
||||
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L${HCC_HOME}/lib -lmcwamp -Wl,-Bsymbolic")
|
||||
add_library(hip_hcc SHARED ${SOURCE_FILES})
|
||||
add_library(hip_hcc SHARED ${SOURCE_FILES_SHARED})
|
||||
add_library(hip_device STATIC ${SOURCE_FILES_STATIC})
|
||||
add_dependencies(hip_device hip_hcc)
|
||||
endif()
|
||||
|
||||
# Generate hcc_version.txt
|
||||
add_custom_target(query_hcc_version COMMAND ${HCC_HOME}/bin/hcc --version > ${PROJECT_BINARY_DIR}/hcc_version.tmp)
|
||||
add_custom_target(check_hcc_version COMMAND ${CMAKE_COMMAND} -E copy_if_different ${PROJECT_BINARY_DIR}/hcc_version.tmp ${PROJECT_BINARY_DIR}/hcc_version.txt DEPENDS query_hcc_version)
|
||||
set_source_files_properties(${PROJECT_BINARY_DIR}/hcc_version.txt PROPERTIES GENERATED TRUE)
|
||||
set_source_files_properties(${SOURCE_FILES} PROPERTIES OBJECT_DEPENDS ${PROJECT_BINARY_DIR}/hcc_version.txt)
|
||||
set_source_files_properties(${SOURCE_FILES_SHARED} ${SOURCE_FILES_STATIC} PROPERTIES OBJECT_DEPENDS ${PROJECT_BINARY_DIR}/hcc_version.txt)
|
||||
add_dependencies(hip_hcc check_hcc_version update_build_and_version_info)
|
||||
|
||||
# Generate .hipInfo
|
||||
@@ -220,8 +225,10 @@ add_custom_target(doc COMMAND HIP_PATH=${CMAKE_CURRENT_SOURCE_DIR} doxygen ${CMA
|
||||
if(HIP_PLATFORM STREQUAL "hcc")
|
||||
if(${HIP_LIB_TYPE} EQUAL 0)
|
||||
install(DIRECTORY ${PROJECT_BINARY_DIR}/CMakeFiles/hip_hcc.dir/src/ DESTINATION lib)
|
||||
else()
|
||||
elseif(${HIP_LIB_TYPE} EQUAL 1)
|
||||
install(TARGETS hip_hcc DESTINATION lib)
|
||||
else()
|
||||
install(TARGETS hip_hcc hip_device DESTINATION lib)
|
||||
endif()
|
||||
install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/src/hip_ir.ll DESTINATION lib)
|
||||
|
||||
|
||||
+1
-1
@@ -328,7 +328,7 @@ if ($needHipHcc) {
|
||||
} elsif ($HIP_LIB_TYPE eq 1) {
|
||||
substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -lhip_hcc " ;
|
||||
} else {
|
||||
substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -Wl,--rpath=$HIP_PATH/lib -lhip_hcc ";
|
||||
substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -Wl,--rpath=$HIP_PATH/lib -lhip_hcc -lhip_device ";
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -588,40 +588,14 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask
|
||||
#define hipGridDim_y (hc_get_num_groups(1))
|
||||
#define hipGridDim_z (hc_get_num_groups(2))
|
||||
|
||||
// loop unrolling
|
||||
__device__ static inline void* memcpy(void* dst, void* src, size_t size)
|
||||
{
|
||||
uint8_t *dstPtr, *srcPtr;
|
||||
dstPtr = (uint8_t*)dst;
|
||||
srcPtr = (uint8_t*)src;
|
||||
for(uint32_t i=0;i<size;i++) {
|
||||
dstPtr[i] = srcPtr[i];
|
||||
}
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
__device__ static inline void* memset(void* ptr, uint8_t val, size_t size)
|
||||
{
|
||||
uint8_t *dstPtr;
|
||||
dstPtr = (uint8_t*)ptr;
|
||||
for(uint32_t i=0;i<size;i++) {
|
||||
dstPtr[i] = val;
|
||||
}
|
||||
return nullptr;
|
||||
}
|
||||
//extern "C" __device__ void* memcpy(void* dst, void* src, size_t size);
|
||||
//extern "C" __device__ void* memset(void* ptr, uint8_t val, size_t size);
|
||||
|
||||
extern "C" __device__ void* __hip_hc_malloc(size_t);
|
||||
extern "C" __device__ void* __hip_hc_free(void *ptr);
|
||||
|
||||
__device__ static inline void* malloc(size_t size)
|
||||
{
|
||||
return __hip_hc_malloc(size);
|
||||
}
|
||||
|
||||
__device__ static inline void* free(void *ptr)
|
||||
{
|
||||
return __hip_hc_free(ptr);
|
||||
}
|
||||
//extern "C" __device__ void* malloc(size_t size);
|
||||
//extern "C" __device__ void* free(void *ptr);
|
||||
|
||||
extern "C" __device__ char4 __hip_hc_add8pk(char4, char4);
|
||||
extern "C" __device__ char4 __hip_hc_sub8pk(char4, char4);
|
||||
|
||||
@@ -7,6 +7,7 @@ elseif(@HIP_LIB_TYPE@ EQUAL 1)
|
||||
install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.a DESTINATION lib)
|
||||
else()
|
||||
install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib)
|
||||
install(FILES @PROJECT_BINARY_DIR@/libhip_device.a DESTINATION lib)
|
||||
endif()
|
||||
install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib)
|
||||
install(FILES @hip_SOURCE_DIR@/src/hip_ir.ll DESTINATION lib)
|
||||
|
||||
@@ -26,6 +26,184 @@ THE SOFTWARE.
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
//=================================================================================================
|
||||
/*
|
||||
Implementation of malloc and free device functions.
|
||||
|
||||
This is the best place to put them because the device
|
||||
global variables need to be initialized at the start.
|
||||
*/
|
||||
|
||||
#define NUM_PAGES_PER_THREAD 16
|
||||
#define SIZE_OF_PAGE 64
|
||||
#define NUM_THREADS_PER_CU 64
|
||||
#define NUM_CUS_PER_GPU 64
|
||||
#define NUM_PAGES NUM_PAGES_PER_THREAD * NUM_THREADS_PER_CU * NUM_CUS_PER_GPU
|
||||
#define SIZE_MALLOC NUM_PAGES * SIZE_OF_PAGE
|
||||
#define SIZE_OF_HEAP SIZE_MALLOC
|
||||
|
||||
size_t g_malloc_heap_size = SIZE_OF_HEAP;
|
||||
|
||||
__attribute__((address_space(1))) char gpuHeap[SIZE_OF_HEAP];
|
||||
__attribute__((address_space(1))) uint32_t gpuFlags[NUM_PAGES];
|
||||
|
||||
__device__ void *__hip_hc_malloc(size_t size)
|
||||
{
|
||||
char *heap = (char*)gpuHeap;
|
||||
if(size > SIZE_OF_HEAP)
|
||||
{
|
||||
return (void*)nullptr;
|
||||
}
|
||||
uint32_t totalThreads = hipBlockDim_x * hipGridDim_x * hipBlockDim_y * hipGridDim_y * hipBlockDim_z * hipGridDim_z;
|
||||
uint32_t currentWorkItem = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x;
|
||||
|
||||
uint32_t numHeapsPerWorkItem = NUM_PAGES / totalThreads;
|
||||
uint32_t heapSizePerWorkItem = SIZE_OF_HEAP / totalThreads;
|
||||
|
||||
uint32_t stride = size / SIZE_OF_PAGE;
|
||||
uint32_t start = numHeapsPerWorkItem * currentWorkItem;
|
||||
|
||||
uint32_t k=0;
|
||||
|
||||
while(gpuFlags[k] > 0)
|
||||
{
|
||||
k++;
|
||||
}
|
||||
|
||||
for(uint32_t i=0;i<stride-1;i++)
|
||||
{
|
||||
gpuFlags[i+start+k] = 1;
|
||||
}
|
||||
|
||||
gpuFlags[start+stride-1+k] = 2;
|
||||
|
||||
void* ptr = (void*)(heap + heapSizePerWorkItem * currentWorkItem + k*SIZE_OF_PAGE);
|
||||
|
||||
return ptr;
|
||||
}
|
||||
|
||||
__device__ void* __hip_hc_free(void *ptr)
|
||||
{
|
||||
if(ptr == nullptr)
|
||||
{
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
uint32_t offsetByte = (uint64_t)ptr - (uint64_t)gpuHeap;
|
||||
uint32_t offsetPage = offsetByte / SIZE_OF_PAGE;
|
||||
|
||||
while(gpuFlags[offsetPage] != 0) {
|
||||
if(gpuFlags[offsetPage] == 2) {
|
||||
gpuFlags[offsetPage] = 0;
|
||||
offsetPage++;
|
||||
break;
|
||||
} else {
|
||||
gpuFlags[offsetPage] = 0;
|
||||
offsetPage++;
|
||||
}
|
||||
}
|
||||
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
__device__ unsigned __hip_ds_bpermute(int index, unsigned src) {
|
||||
return hc::__amdgcn_ds_bpermute(index, src);
|
||||
}
|
||||
|
||||
__device__ float __hip_ds_bpermutef(int index, float src) {
|
||||
return hc::__amdgcn_ds_bpermute(index, src);
|
||||
}
|
||||
|
||||
__device__ unsigned __hip_ds_permute(int index, unsigned src) {
|
||||
return hc::__amdgcn_ds_permute(index, src);
|
||||
}
|
||||
|
||||
__device__ float __hip_ds_permutef(int index, float src) {
|
||||
return hc::__amdgcn_ds_permute(index, src);
|
||||
}
|
||||
|
||||
__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern) {
|
||||
return hc::__amdgcn_ds_swizzle(src, pattern);
|
||||
}
|
||||
|
||||
__device__ float __hip_ds_swizzlef(float src, int pattern) {
|
||||
return hc::__amdgcn_ds_swizzle(src, pattern);
|
||||
}
|
||||
|
||||
__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl) {
|
||||
return hc::__amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl);
|
||||
}
|
||||
|
||||
#define MASK1 0x00ff00ff
|
||||
#define MASK2 0xff00ff00
|
||||
|
||||
__device__ char4 __hip_hc_add8pk(char4 in1, char4 in2) {
|
||||
char4 out;
|
||||
unsigned one1 = in1.val & MASK1;
|
||||
unsigned one2 = in2.val & MASK1;
|
||||
out.val = (one1 + one2) & MASK1;
|
||||
one1 = in1.val & MASK2;
|
||||
one2 = in2.val & MASK2;
|
||||
out.val = out.val | ((one1 + one2) & MASK2);
|
||||
return out;
|
||||
}
|
||||
|
||||
__device__ char4 __hip_hc_sub8pk(char4 in1, char4 in2) {
|
||||
char4 out;
|
||||
unsigned one1 = in1.val & MASK1;
|
||||
unsigned one2 = in2.val & MASK1;
|
||||
out.val = (one1 - one2) & MASK1;
|
||||
one1 = in1.val & MASK2;
|
||||
one2 = in2.val & MASK2;
|
||||
out.val = out.val | ((one1 - one2) & MASK2);
|
||||
return out;
|
||||
}
|
||||
|
||||
__device__ char4 __hip_hc_mul8pk(char4 in1, char4 in2) {
|
||||
char4 out;
|
||||
unsigned one1 = in1.val & MASK1;
|
||||
unsigned one2 = in2.val & MASK1;
|
||||
out.val = (one1 * one2) & MASK1;
|
||||
one1 = in1.val & MASK2;
|
||||
one2 = in2.val & MASK2;
|
||||
out.val = out.val | ((one1 * one2) & MASK2);
|
||||
return out;
|
||||
}
|
||||
|
||||
// loop unrolling
|
||||
__device__ void* memcpy(void* dst, void* src, size_t size)
|
||||
{
|
||||
uint8_t *dstPtr, *srcPtr;
|
||||
dstPtr = (uint8_t*)dst;
|
||||
srcPtr = (uint8_t*)src;
|
||||
for(uint32_t i=0;i<size;i++) {
|
||||
dstPtr[i] = srcPtr[i];
|
||||
}
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
__device__ void* memset(void* ptr, uint8_t val, size_t size)
|
||||
{
|
||||
uint8_t *dstPtr;
|
||||
dstPtr = (uint8_t*)ptr;
|
||||
for(uint32_t i=0;i<size;i++) {
|
||||
dstPtr[i] = val;
|
||||
}
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
__device__ void* malloc(size_t size)
|
||||
{
|
||||
return __hip_hc_malloc(size);
|
||||
}
|
||||
|
||||
__device__ void* free(void *ptr)
|
||||
{
|
||||
return __hip_hc_free(ptr);
|
||||
}
|
||||
|
||||
//=================================================================================================
|
||||
|
||||
// TODO: Choose whether default is precise math or fast math based on compilation flag.
|
||||
#ifdef __HCC_ACCELERATOR__
|
||||
using namespace hc::precise_math;
|
||||
@@ -2294,6 +2472,10 @@ __HIP_DEVICE__ double4 make_double4(double x, double y, double z, double w)
|
||||
return d4;
|
||||
}
|
||||
|
||||
__device__ void __threadfence_system(void){
|
||||
// no-op
|
||||
}
|
||||
|
||||
float __hip_host_erfinvf(float x)
|
||||
{
|
||||
float ret;
|
||||
|
||||
@@ -109,150 +109,6 @@ std::atomic<int> g_lastShortTid(1);
|
||||
std::vector<ProfTrigger> g_dbStartTriggers;
|
||||
std::vector<ProfTrigger> g_dbStopTriggers;
|
||||
|
||||
|
||||
|
||||
/*
|
||||
Implementation of malloc and free device functions.
|
||||
|
||||
This is the best place to put them because the device
|
||||
global variables need to be initialized at the start.
|
||||
*/
|
||||
|
||||
#define NUM_PAGES_PER_THREAD 16
|
||||
#define SIZE_OF_PAGE 64
|
||||
#define NUM_THREADS_PER_CU 64
|
||||
#define NUM_CUS_PER_GPU 64
|
||||
#define NUM_PAGES NUM_PAGES_PER_THREAD * NUM_THREADS_PER_CU * NUM_CUS_PER_GPU
|
||||
#define SIZE_MALLOC NUM_PAGES * SIZE_OF_PAGE
|
||||
#define SIZE_OF_HEAP SIZE_MALLOC
|
||||
|
||||
size_t g_malloc_heap_size = SIZE_OF_HEAP;
|
||||
|
||||
__attribute__((address_space(1))) char gpuHeap[SIZE_OF_HEAP];
|
||||
__attribute__((address_space(1))) uint32_t gpuFlags[NUM_PAGES];
|
||||
|
||||
__device__ void *__hip_hc_malloc(size_t size)
|
||||
{
|
||||
char *heap = (char*)gpuHeap;
|
||||
if(size > SIZE_OF_HEAP)
|
||||
{
|
||||
return (void*)nullptr;
|
||||
}
|
||||
uint32_t totalThreads = hipBlockDim_x * hipGridDim_x * hipBlockDim_y * hipGridDim_y * hipBlockDim_z * hipGridDim_z;
|
||||
uint32_t currentWorkItem = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x;
|
||||
|
||||
uint32_t numHeapsPerWorkItem = NUM_PAGES / totalThreads;
|
||||
uint32_t heapSizePerWorkItem = SIZE_OF_HEAP / totalThreads;
|
||||
|
||||
uint32_t stride = size / SIZE_OF_PAGE;
|
||||
uint32_t start = numHeapsPerWorkItem * currentWorkItem;
|
||||
|
||||
uint32_t k=0;
|
||||
|
||||
while(gpuFlags[k] > 0)
|
||||
{
|
||||
k++;
|
||||
}
|
||||
|
||||
for(uint32_t i=0;i<stride-1;i++)
|
||||
{
|
||||
gpuFlags[i+start+k] = 1;
|
||||
}
|
||||
|
||||
gpuFlags[start+stride-1+k] = 2;
|
||||
|
||||
void* ptr = (void*)(heap + heapSizePerWorkItem * currentWorkItem + k*SIZE_OF_PAGE);
|
||||
|
||||
return ptr;
|
||||
}
|
||||
|
||||
__device__ void* __hip_hc_free(void *ptr)
|
||||
{
|
||||
if(ptr == nullptr)
|
||||
{
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
uint32_t offsetByte = (uint64_t)ptr - (uint64_t)gpuHeap;
|
||||
uint32_t offsetPage = offsetByte / SIZE_OF_PAGE;
|
||||
|
||||
while(gpuFlags[offsetPage] != 0) {
|
||||
if(gpuFlags[offsetPage] == 2) {
|
||||
gpuFlags[offsetPage] = 0;
|
||||
offsetPage++;
|
||||
break;
|
||||
} else {
|
||||
gpuFlags[offsetPage] = 0;
|
||||
offsetPage++;
|
||||
}
|
||||
}
|
||||
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
__device__ unsigned __hip_ds_bpermute(int index, unsigned src) {
|
||||
return hc::__amdgcn_ds_bpermute(index, src);
|
||||
}
|
||||
|
||||
__device__ float __hip_ds_bpermutef(int index, float src) {
|
||||
return hc::__amdgcn_ds_bpermute(index, src);
|
||||
}
|
||||
|
||||
__device__ unsigned __hip_ds_permute(int index, unsigned src) {
|
||||
return hc::__amdgcn_ds_permute(index, src);
|
||||
}
|
||||
|
||||
__device__ float __hip_ds_permutef(int index, float src) {
|
||||
return hc::__amdgcn_ds_permute(index, src);
|
||||
}
|
||||
|
||||
__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern) {
|
||||
return hc::__amdgcn_ds_swizzle(src, pattern);
|
||||
}
|
||||
|
||||
__device__ float __hip_ds_swizzlef(float src, int pattern) {
|
||||
return hc::__amdgcn_ds_swizzle(src, pattern);
|
||||
}
|
||||
|
||||
__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl) {
|
||||
return hc::__amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl);
|
||||
}
|
||||
|
||||
#define MASK1 0x00ff00ff
|
||||
#define MASK2 0xff00ff00
|
||||
|
||||
__device__ char4 __hip_hc_add8pk(char4 in1, char4 in2) {
|
||||
char4 out;
|
||||
unsigned one1 = in1.val & MASK1;
|
||||
unsigned one2 = in2.val & MASK1;
|
||||
out.val = (one1 + one2) & MASK1;
|
||||
one1 = in1.val & MASK2;
|
||||
one2 = in2.val & MASK2;
|
||||
out.val = out.val | ((one1 + one2) & MASK2);
|
||||
return out;
|
||||
}
|
||||
|
||||
__device__ char4 __hip_hc_sub8pk(char4 in1, char4 in2) {
|
||||
char4 out;
|
||||
unsigned one1 = in1.val & MASK1;
|
||||
unsigned one2 = in2.val & MASK1;
|
||||
out.val = (one1 - one2) & MASK1;
|
||||
one1 = in1.val & MASK2;
|
||||
one2 = in2.val & MASK2;
|
||||
out.val = out.val | ((one1 - one2) & MASK2);
|
||||
return out;
|
||||
}
|
||||
|
||||
__device__ char4 __hip_hc_mul8pk(char4 in1, char4 in2) {
|
||||
char4 out;
|
||||
unsigned one1 = in1.val & MASK1;
|
||||
unsigned one2 = in2.val & MASK1;
|
||||
out.val = (one1 * one2) & MASK1;
|
||||
one1 = in1.val & MASK2;
|
||||
one2 = in2.val & MASK2;
|
||||
out.val = out.val | ((one1 * one2) & MASK2);
|
||||
return out;
|
||||
}
|
||||
//=================================================================================================
|
||||
// Thread-local storage:
|
||||
//=================================================================================================
|
||||
|
||||
@@ -1037,7 +1037,3 @@ hipError_t hipFreeArray(hipArray* array)
|
||||
return ihipLogStatus(hipStatus);
|
||||
}
|
||||
|
||||
|
||||
__device__ void __threadfence_system(void){
|
||||
// no-op
|
||||
}
|
||||
|
||||
Ссылка в новой задаче
Block a user