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

Adding checks and debug output for fat binary for hip-clang

Browse Source 
[ROCm/clr commit: da6c5c152e]
This commit is contained in:
Yaxun Sam Liu
2018-10-25 12:19:32 -04:00
parent bf6cc74f25
commit 4012f63d87
1 changed files with 41 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/src/hip_clang.cpp
+41 -12
View File
@@ -86,6 +86,7 @@ __hipRegisterFatBinary(const void* data)
std::string target{&desc->triple[sizeof(AMDGCN_AMDHSA_TRIPLE)],
desc->tripleSize - sizeof(AMDGCN_AMDHSA_TRIPLE)};
tprintf(DB_FB, "Found bundle for %s\n", target.c_str());
for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) {
hsa_agent_t agent = g_allAgents[deviceId + 1];
@@ -110,10 +111,25 @@ __hipRegisterFatBinary(const void* data)
if (module->executable.handle) {
modules->at(deviceId) = module;
tprintf(DB_FB, "Loaded code object for %s\n", name);
} else {
fprintf(stderr, "Failed to load code object for %s\n", name);
abort();
}
}
}
for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) {
hsa_agent_t agent = g_allAgents[deviceId + 1];
char name[64] = {};
hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, name);
if (!(*modules)[deviceId]) {
fprintf(stderr, "No device code bundle for %s\n", name);
abort();
}
}
tprintf(DB_FB, "__hipRegisterFatBinary succeeds and returns %p\n", modules);
return modules;
}
@@ -132,13 +148,18 @@ extern "C" void __hipRegisterFunction(
dim3* gridDim,
int* wSize)
{
HIP_INIT_API(modules, hostFunction, deviceFunction, deviceName);
std::vector<hipFunction_t> functions{g_deviceCnt};
assert(modules && modules->size() >= g_deviceCnt);
for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) {
hipFunction_t function;
if (hipSuccess == hipModuleGetFunction(&function, modules->at(deviceId), deviceName)) {
functions[deviceId] = function;
}
else {
tprintf(DB_FB, "missing kernel %s for device %d\n", deviceName, deviceId);
}
}
g_functions.insert(std::make_pair(hostFunction, std::move(functions)));
@@ -180,6 +201,7 @@ hipError_t hipSetupArgument(
size_t size,
size_t offset)
{
HIP_INIT_API(arg, size, offset);
auto ctx = ihipGetTlsDefaultCtx();
LockedAccessor_CtxCrit_t crit(ctx->criticalData());
auto& arguments = crit->_execStack.top()._arguments;
@@ -194,6 +216,7 @@ hipError_t hipSetupArgument(
hipError_t hipLaunchByPtr(const void *hostFunction)
{
HIP_INIT_API(hostFunction);
ihipExec_t exec;
{
auto ctx = ihipGetTlsDefaultCtx();
@@ -213,20 +236,26 @@ hipError_t hipLaunchByPtr(const void *hostFunction)
deviceId = 0;
}
hipError_t e = hipSuccess;
decltype(g_functions)::iterator it;
if ((it = g_functions.find(hostFunction)) == g_functions.end())
return hipErrorUnknown;
if ((it = g_functions.find(hostFunction)) == g_functions.end()) {
e = hipErrorUnknown;
fprintf(stderr, "kernel %p not found!\n", hostFunction);
abort();
} else {
size_t size = exec._arguments.size();
void *extra[] = {
HIP_LAUNCH_PARAM_BUFFER_POINTER, &exec._arguments[0],
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
HIP_LAUNCH_PARAM_END
};
size_t size = exec._arguments.size();
void *extra[] = {
HIP_LAUNCH_PARAM_BUFFER_POINTER, &exec._arguments[0],
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
HIP_LAUNCH_PARAM_END
};
e = hipModuleLaunchKernel(it->second[deviceId],
exec._gridDim.x, exec._gridDim.y, exec._gridDim.z,
exec._blockDim.x, exec._blockDim.y, exec._blockDim.z,
exec._sharedMem, exec._hStream, nullptr, extra);
}
return hipModuleLaunchKernel(it->second[deviceId],
exec._gridDim.x, exec._gridDim.y, exec._gridDim.z,
exec._blockDim.x, exec._blockDim.y, exec._blockDim.z,
exec._sharedMem, exec._hStream, nullptr, extra);
return ihipLogStatus(e);
}