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

Added hipDeviceGetLimit api

Browse Source 
1. hipDeviceGetLimit API for HCC path is added
2. Test for hipDeviceGetLimit API is added
3. The feature added only supports querying heap size
4. Corrected indents for malloc and free device functions
5. Removed redundant data structures
6. Added g_heap_malloc_size to store the heap size

Change-Id: If48d1b0ce9270e994f1c542cc283ddbb14746bbb


[ROCm/hip commit: 62ec53740c]
This commit is contained in:
Aditya Atluri
2016-10-12 19:58:48 -05:00
parent a11cd21801
commit ef0cb7d619
4 changed files with 80 additions and 45 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/src/hip_hcc.cpp
+34 -40
View File
@@ -86,8 +86,6 @@ hsa_agent_t g_cpu_agent;
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
@@ -98,47 +96,44 @@ hsa_agent_t g_cpu_agent;
#define SIZE_MALLOC NUM_PAGES * SIZE_OF_PAGE
#define SIZE_OF_HEAP SIZE_MALLOC
struct heapTracker_t {
void *ptr;
uint32_t *flags;
uint32_t next;
};
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;
__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 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 stride = size / SIZE_OF_PAGE;
uint32_t start = numHeapsPerWorkItem * currentWorkItem;
uint32_t k=0;
uint32_t k=0;
while(gpuFlags[k] > 0)
{
k++;
}
while(gpuFlags[k] > 0)
{
k++;
}
for(uint32_t i=0;i<stride-1;i++)
{
gpuFlags[i+start+k] = 1;
}
for(uint32_t i=0;i<stride-1;i++)
{
gpuFlags[i+start+k] = 1;
}
gpuFlags[start+stride-1+k] = 2;
gpuFlags[start+stride-1+k] = 2;
void *ptr = (void*)(heap + heapSizePerWorkItem * currentWorkItem + k*SIZE_OF_PAGE);
void* ptr = (void*)(heap + heapSizePerWorkItem * currentWorkItem + k*SIZE_OF_PAGE);
return ptr;
return ptr;
}
__device__ void* __hip_hc_free(void *ptr)
@@ -151,15 +146,14 @@ __device__ void* __hip_hc_free(void *ptr)
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++;
while(gpuFlags[offsetPage] != 0) {
if(gpuFlags[offsetPage] == 2) {
gpuFlags[offsetPage] = 0;
offsetPage++;
break;
} else {
gpuFlags[offsetPage] = 0;
offsetPage++;
}
}