`test/unit/queue/queueTest` fails on Radeon Pro W7800 (`gfx1100`) with ROCm 6.2.4 and later

[fwyzard]

$ ./test/unit/queue/queueTest 
Randomness seeded to: 1375303555
BEGIN not-awaited task in Queue 'alpaka::QueueGenericThreadsNonBlocking<alpaka::DevCpu>' (if there is no matching 'END' line, the task ran past program termination)
END not-awaited task in Queue 'alpaka::QueueGenericThreadsNonBlocking<alpaka::DevCpu>'
BEGIN not-awaited task in Queue 'alpaka::uniform_cuda_hip::detail::QueueUniformCudaHipRt<alpaka::ApiHipRt, false>' (if there is no matching 'END' line, the task ran past program termination)
END not-awaited task in Queue 'alpaka::uniform_cuda_hip::detail::QueueUniformCudaHipRt<alpaka::ApiHipRt, false>'

~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
queueTest is a Catch2 v3.5.2 host application.
Run with -? for options

-------------------------------------------------------------------------------
enqueueBenchmark - alpaka::test::TestQueues - 0
-------------------------------------------------------------------------------
[...]
-------------------------------------------------------------------------------
enqueueBenchmark - alpaka::test::TestQueues - 1
-------------------------------------------------------------------------------
[...]
-------------------------------------------------------------------------------
enqueueBenchmark - alpaka::test::TestQueues - 2
-------------------------------------------------------------------------------
[...]
/data/user/fwyzard/hackathon/alpaka/test/unit/queue/src/QueueTest.cpp:256: FAILED:
due to a fatal error condition:
  SIGSEGV - Segmentation violation signal

===============================================================================
test cases: 35 | 34 passed | 1 failed
assertions: 65 | 64 passed | 1 failed

Segmentation fault (core dumped)

Configured and built with

cmake ../alpaka \
  -Dalpaka_ACC_CPU_B_SEQ_T_SEQ_ENABLE=ON \
  -Dalpaka_ACC_GPU_HIP_ENABLE=ON \
  -DCMAKE_HIP_PLATFORM=amd \
  -DCMAKE_HIP_ARCHITECTURES=gfx1100 \
  -DCMAKE_HIP_COMPILER=/opt/rocm-6.2.4/llvm/bin/clang++ \
  -DCMAKE_PREFIX_PATH=/opt/rocm-6.2.4/lib/cmake \
  -Dalpaka_BUILD_EXAMPLES=ON \
  -Dalpaka_BUILD_BENCHMARKS=ON \
  -DBUILD_TESTING=ON

make

Tested with ROCm 5.7.3, 6.2.4, 6.3.2, 6.4.0, using the AMD driver from 6.3.2 and 6.4.0.
The 5.7 build works, the 6.x builds fail as above.

[fwyzard]

The problem is related to ROCm, not to alpaka.

Here is a native HIP reproducer, that enqueues a host function that does nothing:

// C++ standard headers
#include <cassert>
#include <cstdlib>

// ROCm headers
#include <hip/hip_runtime.h>

#define HIP_CHECK(ARG, ...) assert((ARG) == hipSuccess) 

int main(int argc, char** argv)
{
    int reps = 55000;
    if (argc > 1) {
        reps = std::atoi(argv[1]);
    }

    HIP_CHECK(hipSetDevice(0));

    hipStream_t queue;
    HIP_CHECK(hipStreamCreate(&queue));

    for (int i = 0; i < reps; i++) {
        HIP_CHECK(hipLaunchHostFunc(
            queue,
            [](void*) {},
            nullptr)
        );
    }

    HIP_CHECK(hipStreamSynchronize(queue));
    HIP_CHECK(hipStreamDestroy(queue));  // if reps > ~53000, this calls causes a segmentation fault 
    HIP_CHECK(hipDeviceReset());

}

Using the AMD drivers from ROCm 6.4.0, the reproducer fails with all versions of ROCm at hand (5.7, 6.2, 6.4).

[fwyzard]

Link to the upstream bug report to ROCm: ROCm/ROCm#4751 .

[SimeonEhrig]

I detected the problem months ago on our workstation, that the Radeon RX 7900XTX does not work with ROCm 6.2.4. The W7800 uses the same Chip with RDNA 3 architecture. Unfortunately, I had never time to investigate the problem (or reporting it except to @psychocoderHPC - sorry for that). Thanks for doing it.

[fwyzard]

Did it work for you with an older driver and ROCm version ?
I tried reverting to 5.7, but I still got the same error.

[SimeonEhrig]

On HAL, we use the latest AMDGPU driver:

amdgpu-core/jammy,now 1:6.4.60400-2147987.22.04 all [installed,automatic]
amdgpu-dkms-firmware/jammy,now 1:6.12.12.60400-2147987.22.04 all [installed]
amdgpu-dkms/jammy,now 1:6.12.12.60400-2147987.22.04 all [installed]

and ROCm 6.0.2

[SimeonEhrig]

I tested your reproducer on a Radeon VII (ROCm 6.2.4 and AMDGPU 6.12) and I get the same segmentation fault.