From c0419cc74982c60fcb41614d2bee3631a5ba0988 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Fri, 18 Nov 2016 14:33:20 +0530
Subject: [PATCH] Refactor for building HIP as dynamic library

Change-Id: I65a3d9d589c4fdbbdcf1611e5427224253be8260
---
 CMakeLists.txt                       |  23 ++--
 bin/hipcc                            |   2 +-
 include/hip/hcc_detail/hip_runtime.h |  34 +----
 packaging/hip_hcc.txt                |   1 +
 src/device_util.cpp                  | 182 +++++++++++++++++++++++++++
 src/hip_hcc.cpp                      | 144 ---------------------
 src/hip_memory.cpp                   |   4 -
 7 files changed, 203 insertions(+), 187 deletions(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 885897fb1b..1b95145b96 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -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)
 
diff --git a/bin/hipcc b/bin/hipcc
index dee0894869..641f70b065 100755
--- a/bin/hipcc
+++ b/bin/hipcc
@@ -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 ";
     }
 }
 
diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h
index 44a61e6924..d7c8d3a675 100644
--- a/include/hip/hcc_detail/hip_runtime.h
+++ b/include/hip/hcc_detail/hip_runtime.h
@@ -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);
diff --git a/packaging/hip_hcc.txt b/packaging/hip_hcc.txt
index d7b0877b62..78717247f0 100644
--- a/packaging/hip_hcc.txt
+++ b/packaging/hip_hcc.txt
@@ -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)
diff --git a/src/device_util.cpp b/src/device_util.cpp
index b267c8cbd7..deb0db0a34 100644
--- a/src/device_util.cpp
+++ b/src/device_util.cpp
@@ -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;
diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp
index 3ebe9c3647..c7afd50171 100644
--- a/src/hip_hcc.cpp
+++ b/src/hip_hcc.cpp
@@ -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:
 //=================================================================================================
diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index d3bbd9aeab..089d1a626e 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -1037,7 +1037,3 @@ hipError_t hipFreeArray(hipArray* array)
     return ihipLogStatus(hipStatus);
 }
 
-
-__device__ void  __threadfence_system(void){
-    // no-op
-}