hipFreeAsync hangs
#3370
Comments
|
I also ran tests using debug Julia & HIP build and besides hitting this assert (which I commented out) there were no other issues. |
|
Hi @pxl-th can you please attach a reproducer for the issue. Can you reproduce the hang in C++ as well? |
|
Unfortunately, I was unable to create a MWE as it is unclear to me what causes it. When running tests I get a lot of page faults in Also, reproducing the tests with C++ is not easy, because we have almost 13k tests.
At some point, test workers will become idle and inspecting them with I'm also not sure if this is an issue with Julia or AMDGPU.jl package, we've been successfully running CI on RX6700XT for several months now without issues using ROCm 5.4 - 5.6 and tried other GPUs like MI200. |
|
Also, on Windows there are no issues at all with RX7900XT, it passes all AMDGPU.jl tests without hanging. |
|
@iassiour, not sure if this is expected, but I noticed that async malloc/free vs non-async is ~300x slower (tried on RX6700 XT and RX7900 XT). MWE: #include <hip/hip_runtime.h>
#include <iostream>
using namespace std;
void check(int res) {
if (res != 0) {
std::cerr << "Fail" << std::endl;
}
}
int main(int argc, char* argv[]) {
hipStream_t s;
check(hipStreamCreateWithPriority(&s, 0, 0));
/*
std::cout << "Regular" << std::endl;
for (int i = 1; i < 100000; i++) {
float *x;
check(hipMalloc((void**)&x, 4));
check(hipFree(x));
}
*/
std::cout << "Async" << std::endl;
for (int i = 1; i < 100000; i++) {
float *x;
check(hipMallocAsync((void**)&x, 4, s));
check(hipFreeAsync(x, s));
}
return 0;
} |
|
Hi @pxl-th I think the slowness in async malloc/free is caused by a bug triggered by small (<8 byte) allocations. |
|
Indeed, smaller than 8 bytes allocations are much slower. Thanks! As a note, the reason I've stumbled upon this is that users of AMDGPU.jl reported 20x slower performance than CPU when training ODE to solve MNIST. And it progressively was getting worse as you repeat the task (e.g. run training over and over again). Moving to non-async malloc/free led to 6x improvement in performance and stable compute time. |
|
Hi @pxl-th the fix for < 8-byte allocations has been merged in develop ROCm/clr@2ede1c9 and it should appear in future release. Regarding the 16-byte allocations timing test: Regarding the hang in hipFreeAsync mentioned in the original post, I could not immediately reproduce the issue with 5.7.1 ubuntu 22.04 but with a RX7900XTX. Is there a specific subtest that the workers become idle or it happens in the end of the process? If possible can please attach the logs up to the hanging point. |
|
Thank you for the fix! Regarding Now I'm able to successfully run the test suite, however, it still hangs randomly when running tests and doing some graphics stuff at the same time. I was screencasting at the same time as running tests, but just re-running tests without it worked fine (see CI run just below the failed one).
Usually it hangs at some internal synchronization point. So besides suggesting to run the tests and do some graphics related stuff at the same time I'm not sure how else to reproduce it... But at least now CI passes with Navi 3, so that's an improvement :) |
|
Find the smallest test case, and dump the AMD_LOG_LEVEL=4 for it. |
|
There are tests that reliably trigger the hang. For this we have tests that check that TLS is working properly, where we create streams with different priorities and check that TLS is updated accordingly (that are then destroyed one GC collects them). By default those tests are disabled for Navi 3, so I've uncommented them in
|
|
Reviving this as I have a fairly small MWE that consistently reproduces the issue. Again in Julia as it is much easier to set up the code. MWE: using AMDGPU
function main()
data = rand(Float64, 1024, 1024)
Threads.@threads for i in 1:1000
sum(ROCArray(data))
end
end
main()
And at some point during execution it hangs. Output of Output of |
|
Mixing default and non-default streams in Here's C++ reproducer: #include <hip/hip_runtime.h>
#include <thread>
void fn() {
hipStream_t stream;
hipStreamCreate(&stream);
int n_elements = 1024 * 1024;
int size = n_elements * sizeof(int);
int *a = new int[n_elements];
int *b = new int[n_elements];
int *da, *db;
hipMallocAsync(&da, size, stream);
hipMallocAsync(&db, size, stream);
hipMemcpyHtoDAsync(da, a, size, stream);
hipMemcpyHtoDAsync(db, b, size, stream);
/* hipFreeAsync(da, stream); */ // <--- Works fine.
hipFreeAsync(da, nullptr); // <--- Mixing default stream with non-default causes hang!
hipFreeAsync(db, stream);
hipStreamSynchronize(stream);
hipStreamDestroy(stream);
delete[] a;
delete[] b;
}
void thread_fn() {
for (int i = 0; i < 1000; i++) {
fn();
}
}
int main() {
std::thread t1(thread_fn);
std::thread t2(thread_fn);
std::thread t3(thread_fn);
std::thread t4(thread_fn);
t1.join();
t2.join();
t3.join();
t4.join();
return 0;
} |
|
Kind ping, to see if someone can take a look at the issue. |
|
Testing on ROCm 6.1 with RX 7800 XT, the Julia MWE does no longer hang. However, the C++ reproducer cannot complete. |
|
This might be related to this issue: ROCm/hipFFT#91 |
|
@torrance thanks for the update! This should significantly help with CI in AMDGPU.jl |
|
Indeed - thanks! So this should land in ROCm 6.1.1 right |
|
@luraess It's fixed in future release of ROCm 6.1.2 Thanks! |
Hi, I'm experiencing hangs with
hipFreeAsyncand was wondering what could potentially cause that.From my perspective it looks like some kind of racing condition.
It consistently happens at the end of the test suite when we start to release memory of the device arrays used in the process in AMDGPU.jl which provides AMD GPU programming interface in Julia language.
Just to note, that memory free happens a lot during tests, it just that it hangs at the end.
I made sure that we do not destroy streams or respective context.
Also, freeing arrays uses NULL stream, but for other operations we use other streams.
I started seeing this issues with ROCm 5.6-5.7.1 and using RX7900XT.
Here's gdb output of the process when it hangs:
On ROCm 5.4 it was not observed and the whole test suite ran fine.
If you need any additional info, I'm happy to provide.
The text was updated successfully, but these errors were encountered: