Analyzing problems with HIP port

[markdewing]

Track investigations into the HIP port

Snapshot of the state is in #161

Tests with ROCm 4.1
Main 'hip' executable still fails with the same error

Now working:

• prefixScan_t
• eigenSoA_t

Still failing

• radixSort_t
• gpuVertexFinder_t

Broken differently - testEigenGPUNoFit returns an error rather than hangs

./testEigenGPUNoFit
TEST EIGENVALUES
TEST INVERSE 3x3
TEST INVERSE 4x4
TEST INVERSE 5x5
:0:rocdevice.cpp            :2325: 4804578096 us: Device::callbackQueue aborting with status: 0x1016
Aborted (core dumped)

New failure?

./testEigenGPU
...
testEigenGPU: /mnt/nvme/physics/hep/tracking/patatrack/pixeltrack-standalone/src/hip/test/test_common.h:24: bool isEqualFuzzy(C1, C2, double) [C1 = Eigen::Matrix<double, 2, 1, 0, 2, 1>, C2 = Eigen::Matrix<double, 2, 1, 0, 2, 1>]: Assertion `std::abs(a(i, j) - b(i, j)) < std::min(std::abs(a(i, j)), std::abs(b(i, j))) * epsilon' failed.
Aborted (core dumped)

[markdewing]

Trying to get rocgdb to work to make debugging some of these failures a little easier. Still not working.

Saga here: ROCm/ROCgdb#6

Later update: fixed now. Was installing updated linux kernel incorrectly.

[markdewing]

gpuVertexFinder_t hangs in radix sort.
Looking into the failures in radixSort_t

[markdewing]

One issue with radix sort is in the while loop using ibs as the control variable. At the end of the loop, there is the code

if (threadIdx.x == 0)
  ibs -= sb;
__syncthreads();

Apparently the decrement of ibs does not get propagated to the other threads. The loop repeats until the assert right before this code triggers (assert(c[bin] >= 0))
If a printf statement is added to print out the value of ibs, the radix sort works (for some cases).

This change also causes gpuVertexFinder_t finish its run and not hang.

The radix sort test still issues messages about values being out of order.

[markdewing]

With rocm 4.2

Main hip executable:

:0:rocdevice.cpp            :2533: 1736637763963 us: Device::callbackQueue aborting with error : HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION:  The agent attempted to access memory beyond the largest legal address. code: 0x29
Aborted (core dumped)

Fails:

• radixSort_t - fails as usual
• gpuVertexFinder_t - hangs
• testEigenGPU - matrix comparison fails
• testEigenGPUNoFit

:0:rocdevice.cpp            :2533: 1737277242139 us: Device::callbackQueue aborting with error : HSA_STATUS_ERROR_EXCEPTION: An HSAIL operation resulted in a hardware exception. code: 0x1016

[markdewing]

Kokkos with HIP backend (and rocm 4.2)

Main kokkos executable:

Memory access fault by GPU node-1 (Agent handle: 0x55cd4e1d7350) on address 0x7ff110aef000. Reason: Page not present or supervisor privilege.

Fails

• AtomicPairCounter_t.device - device side assertion "c.m < n" failed
• testEigenGPU - matrix comparison fails

[markdewing]

In radixSort, there is a problem with hipcc in the host code that checks the sorted values for correctness. The shorten function fails to affect the argument, and the program then outputs that the values are out of order. This affects tests where the number of significant bytes used for sorting is less than the number of bytes in the value.

A workaround is to make the shorten function take uint64_t as the input type, and convert values to uint64_t before calling the function.

Issue here:
ROCm/hip#2297

[markdewing]

One of the previous problems was with the ibs variable. Using a printf statement as a workaround produces an inconvenient amount of output. It also works to create another shared variable ibs2 and assign ibs to that.

if (threadIdx.x == 0) {
  ibs -= sb;
  ibs2 = ibs;
}
__syncthreads();

[markdewing]

The code branch that uses shared memory also specifies __launch_bounds__ on the kernel. The kernel appears to not actually be run, but the launchKernel API returns no errors. This happens even the bounds are set the default values (according to the documentation - that would be (1024, 1)). The workaround is to the remove the launch bounds entirely.

With these workarounds, the radixSort_t test passes.

[markdewing]

The gpuVertexFinder_t test hangs in radixSort, around line 152:

    // broadcast
    ibs = size - 1;
    __syncthreads();  // hangs here
    while (__syncthreads_and(ibs > 0)) {

A similar workaround with the ibs variable works. Create a new shared variable ibs31 (doesn't work to reuse ibs2) and assign ibs to it.

    // broadcast
    ibs = size - 1;
    ibs3 = ibs;
    __syncthreads();
    while (__syncthreads_and(ibs > 0)) {

With this workaround, the gpuVertexFinder_t test runs to completion.

[markdewing]

Made a stand-alone reproducer for the problem with the ibs variable

ROCm/hip#2298

Closed the issue with the shorten function in the host code, as the code is incorrect. Will submit a PR to fix it.

[VinInn]

in principle the __syncthreads before the while is not needed as it should sync in __syncthreads_and
please try to add __threadfence() instead of those dummy variable (actually __threadfence_block()
https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#memory-fence-functions

[markdewing]

Testing with ROCm 4.3:

• the ibs3 workaround is not longer needed for gpuVertexFinder
• Using __threadfence(); works instead of using the ibs2 workaround.

[VinInn]

excellent news!
I updated the E4 machine to 4.3 yesterday