From 0ff408a56cef3b1dfcd6b622d180bf4ef22de040 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Thu, 1 Nov 2018 16:20:35 -0400 Subject: [PATCH] Move the global arrays for hip malloc/free from a header into a source file such that there's only an unique copy in an executable and prevent wasting static memory on the host Change-Id: Id5b62766f77809c8d7b47892cb7149c490dcbdb9 --- include/hip/hcc_detail/hip_memory.h | 9 ++++----- src/hip_memory.cpp | 3 +++ 2 files changed, 7 insertions(+), 5 deletions(-) diff --git a/include/hip/hcc_detail/hip_memory.h b/include/hip/hcc_detail/hip_memory.h index 739e488ca3..2c9ec1b7c3 100644 --- a/include/hip/hcc_detail/hip_memory.h +++ b/include/hip/hcc_detail/hip_memory.h @@ -27,6 +27,8 @@ THE SOFTWARE. // HIP heap is implemented as a global array with fixed size. Users may define // __HIP_SIZE_OF_PAGE and __HIP_NUM_PAGES to have a larger heap. +#if __HCC__ || __HIP__ + // Size of page in bytes. #ifndef __HIP_SIZE_OF_PAGE #define __HIP_SIZE_OF_PAGE 64 @@ -39,11 +41,8 @@ THE SOFTWARE. #define __HIP_SIZE_OF_HEAP (__HIP_NUM_PAGES * __HIP_SIZE_OF_PAGE) -#if __HCC__ || __HIP__ - -__attribute__((weak)) __device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP]; -__attribute__((weak)) __device__ - uint32_t __hip_device_page_flag[__HIP_NUM_PAGES]; +extern __device__ char __hip_device_heap[]; +extern __device__ uint32_t __hip_device_page_flag[]; extern "C" inline __device__ void* __hip_malloc(size_t size) { char* heap = (char*)__hip_device_heap; diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index 93ac527826..4ea5b24f43 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -28,6 +28,9 @@ THE SOFTWARE. #include "hip_hcc_internal.h" #include "trace_helper.h" +__device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP]; +__device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES]; + // Internal HIP APIS: namespace hip_internal {