AI/rocm-systems
2
0
Fork 0
Koodi Ongelmat Pull-pyynnöt Actions Paketit Projektit Julkaisut Wiki Toiminta

Refactor Module and Function APIs.

Selaa lähdekoodia 
- 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).
This commit is contained in:
Ben Sander
2016-12-17 07:19:22 -06:00
vanhempi 263a9614ff
commit 4d29885be3
4 muutettua tiedostoa jossa 52 lisäystä ja 65 poistoa
Download Patch File Download Diff File Expand all files Collapse all files
src/hip_module.cpp
+35 -29
Näytä tiedosto
@@ -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<void*>(func->_kernel);
*dptr = reinterpret_cast<void*>(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);
}