AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Move the global arrays for hip malloc/free

Browse Source 
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
This commit is contained in:
Siu Chi Chan
2018-11-01 16:20:35 -04:00
parent df2dd762c2
commit 0ff408a56c
2 changed files with 7 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/hip/hcc_detail/hip_memory.h
+4 -5
View File
@@ -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;