From 62ec53740cd383e9cbcf9672c7ce3900296a7ec7 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Wed, 12 Oct 2016 19:58:48 -0500 Subject: [PATCH] Added hipDeviceGetLimit api 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 --- include/hip/hcc_detail/hip_runtime_api.h | 15 +++++ src/hip_device.cpp | 26 ++++++-- src/hip_hcc.cpp | 74 ++++++++++------------ tests/src/deviceLib/hipTestDeviceLimit.cpp | 10 +++ 4 files changed, 80 insertions(+), 45 deletions(-) create mode 100644 tests/src/deviceLib/hipTestDeviceLimit.cpp diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 4cbaf9ea4e..135e307b6e 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -60,6 +60,7 @@ typedef void* hipDeviceptr_t; typedef struct ihipEvent_t *hipEvent_t; +typedef unsigned hipLimit; /** * @addtogroup GlobalDefs More @@ -97,6 +98,8 @@ typedef struct ihipEvent_t *hipEvent_t; #define hipDeviceMapHost 0x8 #define hipDeviceLmemResizeToMax 0x16 +#define hipLimitMallocHeapSize 0x2 + /** * @warning On AMD devices and recent Nvidia devices, these hints and controls are ignored. */ @@ -323,6 +326,18 @@ hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig ); */ hipError_t hipDeviceGetCacheConfig ( hipFuncCache *cacheConfig ); +/** + * @brief Get Resource limits of current device + * + * @param [out] pValue + * @param [in] limit + * + * @returns #hipSuccess, #hipErrorUnsupportedLimit, #hipErrorInvalidValue + * Note: Currently, only hipLimitMallocHeapSize is available + * + */ +hipError_t hipDeviceGetLimit(size_t *pValue, hipLimit limit); + /** * @brief Set Cache configuration for a specific function diff --git a/src/hip_device.cpp b/src/hip_device.cpp index a677402b69..db04985cf1 100644 --- a/src/hip_device.cpp +++ b/src/hip_device.cpp @@ -68,7 +68,7 @@ hipError_t hipGetDeviceCount(int *count) return e; } -hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig ) +hipError_t hipDeviceSetCacheConfig(hipFuncCache cacheConfig) { HIP_INIT_API(cacheConfig); @@ -77,7 +77,7 @@ hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig ) return ihipLogStatus(hipSuccess); } -hipError_t hipDeviceGetCacheConfig ( hipFuncCache *cacheConfig ) +hipError_t hipDeviceGetCacheConfig(hipFuncCache *cacheConfig) { HIP_INIT_API(cacheConfig); @@ -86,7 +86,23 @@ hipError_t hipDeviceGetCacheConfig ( hipFuncCache *cacheConfig ) return ihipLogStatus(hipSuccess); } -hipError_t hipFuncSetCacheConfig ( hipFuncCache cacheConfig ) +extern "C" size_t g_malloc_heap_size; + +hipError_t hipDeviceGetLimit (size_t *pValue, hipLimit limit) +{ + HIP_INIT_API(pValue, limit); + if(pValue == nullptr) { + return ihipLogStatus(hipErrorInvalidValue); + } + if(limit == hipLimitMallocHeapSize) { + *pValue = g_malloc_heap_size; + return ihipLogStatus(hipSuccess); + }else{ + return ihipLogStatus(hipErrorUnsupportedLimit); + } +} + +hipError_t hipFuncSetCacheConfig (hipFuncCache cacheConfig) { HIP_INIT_API(cacheConfig); @@ -95,7 +111,7 @@ hipError_t hipFuncSetCacheConfig ( hipFuncCache cacheConfig ) return ihipLogStatus(hipSuccess); } -hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ) +hipError_t hipDeviceSetSharedMemConfig (hipSharedMemConfig config) { HIP_INIT_API(config); @@ -104,7 +120,7 @@ hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ) return ihipLogStatus(hipSuccess); } -hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig ) +hipError_t hipDeviceGetSharedMemConfig (hipSharedMemConfig *pConfig) { HIP_INIT_API(pConfig); diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index 06f06a65b3..0ba98b8a4c 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -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 +#include +#include + +int main() +{ + size_t heap; + assert(hipSuccess == hipDeviceGetLimit(&heap, hipLimitMallocHeapSize)); + assert(heap == 4194304); +}