23 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_MEMORY_H 24 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_MEMORY_H 30 #if (__HCC__ || __HIP__) && __HIP_ENABLE_DEVICE_MALLOC__ 33 #ifndef __HIP_SIZE_OF_PAGE 34 #define __HIP_SIZE_OF_PAGE 64 38 #ifndef __HIP_NUM_PAGES 39 #define __HIP_NUM_PAGES (16 * 64 * 64) 42 #define __HIP_SIZE_OF_HEAP (__HIP_NUM_PAGES * __HIP_SIZE_OF_PAGE) 44 #if __HIP__ && __HIP_DEVICE_COMPILE__ 45 __attribute__((weak)) __device__
char __hip_device_heap[__HIP_SIZE_OF_HEAP];
46 __attribute__((weak)) __device__
47 uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
49 extern __device__
char __hip_device_heap[];
50 extern __device__ uint32_t __hip_device_page_flag[];
53 extern "C" inline __device__
void* __hip_malloc(
size_t size) {
54 char* heap = (
char*)__hip_device_heap;
55 if (size > __HIP_SIZE_OF_HEAP) {
56 return (
void*)
nullptr;
58 uint32_t totalThreads =
59 hipBlockDim_x * hipGridDim_x * hipBlockDim_y
60 * hipGridDim_y * hipBlockDim_z * hipGridDim_z;
61 uint32_t currentWorkItem = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x
62 + (hipThreadIdx_y + hipBlockDim_y * hipBlockIdx_y) * hipBlockDim_x
63 + (hipThreadIdx_z + hipBlockDim_z * hipBlockIdx_z) * hipBlockDim_x
66 uint32_t numHeapsPerWorkItem = __HIP_NUM_PAGES / totalThreads;
67 uint32_t heapSizePerWorkItem = __HIP_SIZE_OF_HEAP / totalThreads;
69 uint32_t stride = size / __HIP_SIZE_OF_PAGE;
70 uint32_t start = numHeapsPerWorkItem * currentWorkItem;
74 while (__hip_device_page_flag[k] > 0) {
78 for (uint32_t i = 0; i < stride - 1; i++) {
79 __hip_device_page_flag[i + start + k] = 1;
82 __hip_device_page_flag[start + stride - 1 + k] = 2;
84 void* ptr = (
void*)(heap
85 + heapSizePerWorkItem * currentWorkItem + k * __HIP_SIZE_OF_PAGE);
90 extern "C" inline __device__
void* __hip_free(
void* ptr) {
95 uint32_t offsetByte = (uint64_t)ptr - (uint64_t)__hip_device_heap;
96 uint32_t offsetPage = offsetByte / __HIP_SIZE_OF_PAGE;
98 while (__hip_device_page_flag[offsetPage] != 0) {
99 if (__hip_device_page_flag[offsetPage] == 2) {
100 __hip_device_page_flag[offsetPage] = 0;
104 __hip_device_page_flag[offsetPage] = 0;
114 #endif // HIP_INCLUDE_HIP_HCC_DETAIL_HIP_MEMORY_H