From 4012f63d87f20c3b71d8f2888ce0b77615b718a2 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Thu, 25 Oct 2018 12:19:32 -0400 Subject: [PATCH 1/2] Adding checks and debug output for fat binary for hip-clang [ROCm/clr commit: da6c5c152e879fcaa53118ce98f33e578d40760e] --- projects/clr/hipamd/src/hip_clang.cpp | 53 +++++++++++++++++++++------ 1 file changed, 41 insertions(+), 12 deletions(-) diff --git a/projects/clr/hipamd/src/hip_clang.cpp b/projects/clr/hipamd/src/hip_clang.cpp index 15a96d298a..6067edd289 100644 --- a/projects/clr/hipamd/src/hip_clang.cpp +++ b/projects/clr/hipamd/src/hip_clang.cpp @@ -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 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); } From cc101b22eddf8c08c9047b43c29aaa805ed8c5dd Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Fri, 26 Oct 2018 14:11:18 -0400 Subject: [PATCH 2/2] Add HIP_DUMP_CODE_OBJECT [ROCm/clr commit: 0027a54d8bb6b7229b9fbcd050ff438d0e1b4380] --- projects/clr/hipamd/src/hip_clang.cpp | 11 +++++++++++ projects/clr/hipamd/src/hip_hcc.cpp | 6 ++++++ projects/clr/hipamd/src/hip_hcc_internal.h | 2 +- 3 files changed, 18 insertions(+), 1 deletion(-) diff --git a/projects/clr/hipamd/src/hip_clang.cpp b/projects/clr/hipamd/src/hip_clang.cpp index 6067edd289..cfd75df562 100644 --- a/projects/clr/hipamd/src/hip_clang.cpp +++ b/projects/clr/hipamd/src/hip_clang.cpp @@ -22,6 +22,7 @@ THE SOFTWARE. #include #include +#include #include "hip/hip_runtime.h" #include "hip_hcc_internal.h" @@ -112,6 +113,16 @@ __hipRegisterFatBinary(const void* data) if (module->executable.handle) { modules->at(deviceId) = module; tprintf(DB_FB, "Loaded code object for %s\n", name); + if (HIP_DUMP_CODE_OBJECT) { + char fname[30]; + static std::atomic index; + sprintf(fname, "__hip_dump_code_object%04d.o", index++); + tprintf(DB_FB, "Dump code object %s\n", fname); + std::ofstream ofs; + ofs.open(fname, std::ios::binary); + ofs << image; + ofs.close(); + } } else { fprintf(stderr, "Failed to load code object for %s\n", name); abort(); diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index eff93da847..e152e7ba69 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -97,6 +97,8 @@ int HIP_INIT_ALLOC = -1; int HIP_SYNC_STREAM_WAIT = 0; int HIP_FORCE_NULL_STREAM = 0; +int HIP_DUMP_CODE_OBJECT = 0; + #if (__hcc_workweek__ >= 17300) // Make sure we have required bug fix in HCC @@ -1294,6 +1296,10 @@ void HipReadEnv() { "overridden by specifying hipEventReleaseToSystem or hipEventReleaseToDevice flag " "when creating the event."); + READ_ENV_I(release, HIP_DUMP_CODE_OBJECT, 0, + "If set, dump code object as __hip_dump_code_object[nnnn].o in the current directory," + "where nnnn is the index number."); + // Some flags have both compile-time and runtime flags - generate a warning if user enables the // runtime flag but the compile-time flag is disabled. if (HIP_DB && !COMPILE_HIP_DB) { diff --git a/projects/clr/hipamd/src/hip_hcc_internal.h b/projects/clr/hipamd/src/hip_hcc_internal.h index d64a4a4cbe..8102f066de 100644 --- a/projects/clr/hipamd/src/hip_hcc_internal.h +++ b/projects/clr/hipamd/src/hip_hcc_internal.h @@ -83,11 +83,11 @@ extern int HIP_SYNC_NULL_STREAM; extern int HIP_INIT_ALLOC; extern int HIP_FORCE_NULL_STREAM; +extern int HIP_DUMP_CODE_OBJECT; // TODO - remove when this is standard behavior. extern int HCC_OPT_FLUSH; - // Class to assign a short TID to each new thread, for HIP debugging purposes. class TidInfo { public: