From 4d29885be3652c75c75d5d31d7009dc65d0a9815 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Sat, 17 Dec 2016 07:19:22 -0600 Subject: [PATCH] Refactor Module and Function APIs. - hipFunction_t is now returned by value. This eliminates dynamic allocation / memory management complexity in the module. Removed the kernel name so the structure is just 16 bytes now. - Moved the hsa_executable_load_module and hsa_executable_freeze calls to the hipModuleLoad and hipModuleLoadData calls. - Apply sharedMemBytes in hipModuleLaunchKernel to group segment size (not private). --- include/hip/hcc_detail/hip_runtime_api.h | 9 +++- src/hip_hcc.h | 35 +------------ src/hip_module.cpp | 64 +++++++++++++----------- src/trace_helper.h | 9 ++++ 4 files changed, 52 insertions(+), 65 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index b2cfa59d71..e38f1d6982 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -33,7 +33,7 @@ THE SOFTWARE. #include #include -//#include "hip/hip_hcc.h" +#include #if defined (__HCC__) && (__hcc_workweek__ < 16155) #error("This version of HIP requires a newer version of HCC."); @@ -72,7 +72,12 @@ typedef struct ihipIpcEventHandle_t *hipIpcEventHandle_t; typedef struct ihipModule_t *hipModule_t; -typedef struct ihipFunction_t *hipFunction_t; +struct ihipModuleSymbol_t{ + hsa_executable_symbol_t _hsaSymbol; + uint64_t _object; // The kernel object. +}; + +typedef struct ihipModuleSymbol_t hipFunction_t; typedef void* hipDeviceptr_t; diff --git a/src/hip_hcc.h b/src/hip_hcc.h index b5417080e0..4075c2b068 100644 --- a/src/hip_hcc.h +++ b/src/hip_hcc.h @@ -383,26 +383,6 @@ public: size_t psize; }; -class ihipFunction_t{ -public: - ihipFunction_t(const char *name) { - size_t nameSz = strlen(name); - char *kernelName = (char*)malloc(nameSz); - strncpy(kernelName, name, nameSz); - _kernelName = kernelName; - }; - - ~ihipFunction_t() { - if (_kernelName) { - free((void*)_kernelName); - _kernelName = NULL; - }; - }; -public: - const char *_kernelName; - hsa_executable_symbol_t _kernelSymbol; - uint64_t _kernel; -}; class ihipModule_t { public: @@ -412,20 +392,7 @@ public: void *ptr; size_t size; - ihipModule_t() : executable(), object(), fileName(), ptr(nullptr), size(0), hipFunctionTable() {} - ~ihipModule_t() { - for (int i = 0; i < hipFunctionTable.size(); ++i) { - ihipFunction_t *func = hipFunctionTable[i]; - delete func; - } - hipFunctionTable.clear(); - } - - void registerFunction(ihipFunction_t* func) { - hipFunctionTable.push_back(func); - } -private: - std::vector hipFunctionTable; + ihipModule_t() : executable(), object(), fileName(), ptr(nullptr), size(0) {} }; template diff --git a/src/hip_module.cpp b/src/hip_module.cpp index badeac2dcd..cb56fb3b4a 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -35,6 +35,11 @@ THE SOFTWARE. //TODO Use Pool APIs from HCC to get memory regions. +#define CHECK_HSA(hsaStatus, hipStatus) \ +if (hsaStatus != HSA_STATUS_SUCCESS) {\ + return ihipLogStatus(hipStatus);\ +} + namespace hipdrv { hsa_status_t findSystemRegions(hsa_region_t region, void *data){ @@ -154,9 +159,13 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ } status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable); - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorNotInitialized); - } + CHECK_HSA(status, hipErrorNotInitialized); + + status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL); + CHECK_HSA(status, hipErrorNotInitialized); + + status = hsa_executable_freeze((*module)->executable, NULL); + CHECK_HSA(status, hipErrorNotInitialized); } } @@ -189,7 +198,7 @@ hipError_t hipModuleUnload(hipModule_t hmod) return ihipLogStatus(ret); } -hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const char *name){ +hipError_t ihipModuleGetSymbol(hipFunction_t *func, hipModule_t hmod, const char *name){ auto ctx = ihipGetTlsDefaultCtx(); hipError_t ret = hipSuccess; @@ -201,27 +210,21 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch ret = hipErrorInvalidContext; }else{ - *func = new ihipFunction_t(name); - hmod->registerFunction(*func); int deviceId = ctx->getDevice()->_deviceId; ihipDevice_t *currentDevice = ihipGetDevice(deviceId); hsa_agent_t gpuAgent = (hsa_agent_t)currentDevice->_hsaAgent; - hsa_status_t status; - status = hsa_executable_load_code_object(hmod->executable, gpuAgent, hmod->object, NULL); - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorNotInitialized); - } - status = hsa_executable_freeze(hmod->executable, NULL); - status = hsa_executable_get_symbol(hmod->executable, NULL, name, gpuAgent, 0, &(*func)->_kernelSymbol); + + hsa_status_t status; + status = hsa_executable_get_symbol(hmod->executable, NULL, name, gpuAgent, 0, &(func->_hsaSymbol)); if(status != HSA_STATUS_SUCCESS){ return ihipLogStatus(hipErrorNotFound); } - status = hsa_executable_symbol_get_info((*func)->_kernelSymbol, + status = hsa_executable_symbol_get_info(func->_hsaSymbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, - &(*func)->_kernel); + &func->_object); if(status != HSA_STATUS_SUCCESS){ return ihipLogStatus(hipErrorNotFound); @@ -234,7 +237,7 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod, const char *name){ HIP_INIT_API(hfunc, hmod, name); - return ihipModuleGetFunction(hfunc, hmod, name); + return ihipModuleGetSymbol(hfunc, hmod, name); } @@ -275,7 +278,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, } uint32_t groupSegmentSize; - hsa_status_t status = hsa_executable_symbol_get_info(f->_kernelSymbol, + hsa_status_t status = hsa_executable_symbol_get_info(f._hsaSymbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, &groupSegmentSize); if(status != HSA_STATUS_SUCCESS){ @@ -283,20 +286,19 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, } uint32_t privateSegmentSize; - status = hsa_executable_symbol_get_info(f->_kernelSymbol, + status = hsa_executable_symbol_get_info(f._hsaSymbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, &privateSegmentSize); if(status != HSA_STATUS_SUCCESS){ return ihipLogStatus(hipErrorNotFound); } - privateSegmentSize += sharedMemBytes; /* Kernel argument preparation. */ grid_launch_parm lp; - hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, f->_kernelName); + hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, "ModuleKernel"); #if USE_DISPATCH_HSA_KERNEL @@ -315,7 +317,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, aql.grid_size_z = blockDimZ * gridDimZ; aql.group_segment_size = groupSegmentSize; aql.private_segment_size = privateSegmentSize; - aql.kernel_object = f->_kernel; + aql.kernel_object = f._object; aql.setup = 3 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; aql.header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | (1 << HSA_PACKET_HEADER_BARRIER) | @@ -356,7 +358,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, Launch AQL packet */ hStream->launchModuleKernel(*lp.av, signal, blockDimX, blockDimY, blockDimZ, - gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, kern, kernArgSize, f->_kernel); + gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, kern, kernArgSize, f->_object); /* @@ -384,7 +386,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, #endif // USE_DISPATCH_HSA_KERNEL - ihipPostLaunchKernel(f->_kernelName, hStream, lp); + ihipPostLaunchKernel("ModuleKernel", hStream, lp); } @@ -405,9 +407,9 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t *dptr, size_t *bytes, } else{ hipFunction_t func; - ihipModuleGetFunction(&func, hmod, name); + ihipModuleGetSymbol(&func, hmod, name); *bytes = PrintSymbolSizes(hmod->ptr, name) + sizeof(amd_kernel_code_t); - *dptr = reinterpret_cast(func->_kernel); + *dptr = reinterpret_cast(func._object); return ihipLogStatus(ret); } } @@ -419,7 +421,7 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image) hipError_t ret = hipSuccess; if(image == NULL || module == NULL){ return ihipLogStatus(hipErrorNotInitialized); - }else{ + } else { auto ctx = ihipGetTlsDefaultCtx(); *module = new ihipModule_t; int deviceId = ctx->getDevice()->_deviceId; @@ -452,9 +454,13 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image) } status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable); - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorNotInitialized); - } + CHECK_HSA(status, hipErrorNotInitialized); + + status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL); + CHECK_HSA(status, hipErrorNotInitialized); + + status = hsa_executable_freeze((*module)->executable, NULL); + CHECK_HSA(status, hipErrorNotInitialized); } return ihipLogStatus(ret); } diff --git a/src/trace_helper.h b/src/trace_helper.h index e75b492e0c..bde40d0690 100644 --- a/src/trace_helper.h +++ b/src/trace_helper.h @@ -72,6 +72,15 @@ inline std::string ToString(hipEvent_t v) return ss.str(); }; +// hipEvent_t specialization. TODO - maybe add an event ID for debug? +template <> +inline std::string ToString(hipFunction_t v) +{ + std::ostringstream ss; + ss << "0x" << std::hex << v._object; + return ss.str(); +}; + // hipStream_t