mxz297
commented
2 years ago I recompiled rocprofiler from source for both rocm-4.3.1 and rocm-4.5.2 to have a better understanding of the problem. On surface, the infinite recursion happens because the interceptor function for code object freeze operation is calling itself (https://github.com/ROCm-Developer-Tools/rocprofiler/blob/rocm-4.3.x/src/util/hsa_rsrc_factory.cpp#L779).
Previously for rocm-4.3.1, I was able to work around this problem by disabling code object tracking. This workaround no longer works in rocm-4.5.2 as code object tracking is always enabled (https://github.com/ROCm-Developer-Tools/rocprofiler/blob/rocm-4.5.2/src/core/rocprofiler.cpp#L429). A user can still set the field of disabling code object tracking, but then be ignored. It would be really helpful to document these important internal changes.
Now back to the real problem that why the infinite recursion happened:
hsa_status_t HsaRsrcFactory::hsa_executable_freeze_interceptor(hsa_executable_t executable, const char *options) {
std::lock_guard<mutex_t> lck(mutex_);
if (symbols_map_ == NULL) symbols_map_ = new symbols_map_t;
hsa_status_t status = hsa_api_.hsa_executable_iterate_symbols(executable, executable_symbols_cb, NULL);
CHECK_STATUS("Error in iterating executable symbols", status);
return hsa_api_.hsa_executable_freeze(executable, options);
}The function pointer in the return statement ends up with being HsaRsrcFactory::hsa_executable_freeze_interceptor causing the recursion. hsa_api_ is set in function HsaRsrcFactory::InitHsaApiTable (https://github.com/ROCm-Developer-Tools/rocprofiler/blob/rocm-4.3.x/src/util/hsa_rsrc_factory.cpp#L184). Based on my understanding, hsa_api_ records a set of actual HSA calls, so the interceptors can be interposed upon HSA calls.
HsaRsrcFactory::InitHsaApiTable is called two places:
- https://github.com/ROCm-Developer-Tools/rocprofiler/blob/rocm-4.3.x/src/util/hsa_rsrc_factory.cpp#L127
- https://github.com/ROCm-Developer-Tools/rocprofiler/blob/rocm-4.3.x/src/core/rocprofiler.cpp#L114
When I trace in gdb, the following event happened:
InitHsaApiTableis called in the first call site with input parameterNULL. This leads tohsa_api_to be initialized with function pointers to HSA API entriesInitHsaApiTableis called in the second call site with an input parameter, which represents the actual implementation of HSA APIs. This call does not update tohsa_api_due to the first if statement at the beginning ofInitHsaApiTable.
The fundamental issue here is that HSA API entries are just a wrapper function around the actual implementation function. For example, the API entry for hsa_executable_freeze is shown as follow (I do not find HSA runtime source code, so I just disassemble the shared library)
0000000000067940 <hsa_executable_freeze>:
67940: 48 8b 05 89 2a 42 00 mov 0x422a89(%rip),%rax # 48a3d0 <_ZL12coreApiTable>
67947: ff a0 c0 02 00 00 jmpq *0x2c0(%rax)It is clear to me that this function just reads a function table and then do a tail call to the actual implementation.
Now the problem is that if hsa_api_ points to this public entry function, which is just a wrapper, and then later rocprofiler updates the actual coreAPITable with its interceptor, we end up with the interceptor calling itself.
To resolve this problem, it looks to me that we should just remove the first if state in InitHsaApiTable (https://github.com/ROCm-Developer-Tools/rocprofiler/blob/rocm-4.3.x/src/util/hsa_rsrc_factory.cpp#L187). This ensures that rocprofiler can get the actual hsa implementation functions provided by HSA runtime and can call hsa function when the hsa implementation table is not ready. With this change, at least locally I can resolve the infinite recursion problem.
While I would like to make a PR for this, I find that the rocprofiler github repo is in a quite strange state: Tag rocm-4.5.2 is shown to contain commits not in the repo. Both amd-master branch and the rocm-4.5.2 branch are behind the rocm-4.5.2 tag. Against which branch should I make the PR?
kikimych
marklawsAMD
ppanchad-amd
sohaibnd
Hi,
I am a developer from the HPCToolkit project at Rice University. I am developing AMD GPU counter support in HPCToolkit directly using rocprofiler API. I am currently running into an infinite recursion in librocprofiler.so with the following stack trace:
You can see this is with rocm-4.3.1. With rocm-4.3.1, I was able to work around this issue by
settings->code_obj_tracking = 0;insideOnLoadToolProp:https://github.com/HPCToolkit/hpctoolkit/blob/rocprofiler_support/src/tool/hpcrun/gpu/amd/rocprofiler-api.c#L436
This work around does not seem to work with rocm-4.5.0 and I am seeing a similar infinite recursion in librocprofiler.
Are there any recommendations or insights on resolving this problem?