[AMDGPU] Miscompiled post-barrier LDS reload in a large-workgroup HIP kernel on gfx11 (silent wrong results)

# [AMDGPU] Miscompiled post-barrier LDS reload in a large-workgroup HIP kernel on gfx11 (silent wrong results)

## Summary

On `gfx1100` (RDNA3), a value written to LDS (`__shared__`) by one lane, published
with `__syncthreads()`, and then reloaded by the other lanes of the workgroup can
be read back **inconsistently across lanes** when the workgroup exceeds ~512
threads. The reload appears to be miscompiled: different lanes proceed with
different values of a quantity that must be uniform, producing silently wrong
numerical results (no fault, no assertion).

This was found in rocSOLVER's small-matrix Jacobi eigensolver (`run_syevj`), where
a Jacobi rotation `(c, s)` is computed once by lane `tiy==0`, stored in LDS, and
reloaded by every row-parallel lane after a barrier to apply the rotation. Above
~512 threads the reloaded `(c, s)` differ between lanes, so different matrix rows
receive different rotations and the accumulated eigenvectors lose orthogonality.

## Environment

- GPU: AMD Radeon RX 7900 XT, `gfx1100` (RDNA3, wave32)
- clang / LLVM: `clang version 21.1.8`
- HIP: `7.1.52801`
- Compile: `hipcc -O3 --offload-arch=gfx1100 --gpu-max-threads-per-block=1024`
- Reproduces at both `-O2` and `-O3`.

## Reproducer

Self-contained, ~180 lines, depends only on `hip_runtime` (attached:
`llvm_repro_syevj.cpp`). It runs a faithful transcription of the rocSOLVER kernel
on a batch of ill-conditioned symmetric matrices and checks that the returned
eigenvector matrix `A` is orthogonal (`||A^T A - I|| ~ 1e-13`).

A single `-D` toggles only how the rotation `(c, s)` reaches the applying lanes:

- `BUG=1` (default): compute in lane `tiy==0`, store to `cosines_res/sines_diag`
in LDS, `__syncthreads()`, reload in every lane. ← upstream form
- `BUG=0`: compute `(c, s)` in every lane (a pure function of the same, unmodified
`Acpy` entries, so bit-identical across lanes). No LDS round-trip.

Both are numerically identical; the apply code that follows is byte-for-byte the
same. Only the LDS reload differs.

```
$ hipcc -O3 --offload-arch=gfx1100 --gpu-max-threads-per-block=1024 -DBUG=1 llvm_repro_syevj.cpp -o r_bug && ./r_bug
device=AMD Radeon RX 7900 XT arch=gfx1100 BUG=1
n=39 threads= 400 [small] ||A^T A - I|| worst=6.311e-14 corrupt=0/12 ok
n=48 threads= 576 [small] ||A^T A - I|| worst=4.389e+00 corrupt=12/12 <== MISCOMPILE
n=56 threads= 784 [small] ||A^T A - I|| worst=3.589e+00 corrupt=9/12 <== MISCOMPILE
n=58 threads= 841 [small] ||A^T A - I|| worst=1.047e-13 corrupt=0/12 ok
REPRODUCED (non-orthogonal eigenvectors)

$ hipcc -O3 --offload-arch=gfx1100 --gpu-max-threads-per-block=1024 -DBUG=0 llvm_repro_syevj.cpp -o r_fix && ./r_fix
... all rows: ||A^T A - I|| ~ 1e-13, clean
```

Notes:
- `n=39` → 400 threads (< 512): clean. The corruption only appears once the
workgroup crosses ~512 threads (16 wave32 waves), independent of the matrix size
(the algorithm itself is correct — `BUG=0` is clean at the same sizes).
- The failure is data dependent and, for thread counts that do not pack into whole
wavefronts (e.g. n=58 → 841 threads), non-deterministic run to run. `n=48`
(576 threads = 18 whole waves) reproduces deterministically at 12/12; a given
seed may or may not trip `n=58`. Run a few times / vary the seed to see the
intermittent sizes.

## What was ruled out (rocSOLVER-side investigation)

- Not a memory-model / visibility race: inserting `__threadfence()` before every
barrier is a no-op (bit-identical output).
- Not compiler load/store reordering across the barrier at source level: an
`asm volatile("" ::: "memory")` clobber before the apply has no effect.
- Not a broken `s_barrier`: launching extra no-op lanes past 512 threads (which
hit the barriers but do no writes) stays clean.
- Not register spilling: the kernel does not spill (VGPRs 42, ScratchSize 0).
- Not wave-size selection: rebuilding the kernel `-mwavefrontsize64` does not fix
it (if anything, more failures).
- Not `-O3`-specific: reproduces at `-O2` as well.
- Not inlining scope: moving the apply into a `__noinline__` device function does
not fix it.
- The identical row-parallel global read-modify-write pattern, in isolation
(a small standalone kernel without the surrounding convergence/norm code), is
bit-exact at 841 threads — so the write pattern and the hardware are fine. The
defect only appears in the full kernel context, and is removed by eliminating
the post-barrier LDS reload.

Taken together this points to a codegen defect around the LDS reload of a
barrier-published value under this kernel's control flow / occupancy on gfx11,
rather than an application bug (the two forms are numerically equivalent).

## Ask

- Confirmation of the miscompile and, if possible, identification of the pass /
condition (the ~512-thread threshold and whole-wave-vs-partial-wave determinism
split suggest a scheduling/occupancy interaction).
- Happy to provide LLVM IR / `-save-temps`, an `llvm-reduce`-minimized case, or
ISA on request.

## Downstream

Fixed downstream in rocSOLVER by computing the rotation per-lane instead of
broadcasting it through LDS (equivalent to `BUG=0` above): ROCm/rocm-libraries#9135. That restores correct, deterministic results at no
measurable cost.

Reproducer — llvm_repro_syevj.cpp

```cpp
// Self-contained reproducer for an LLVM AMDGPU miscompile on gfx11 (RDNA3).
//
// hipcc -O3 --offload-arch=gfx1100 --gpu-max-threads-per-block=1024 llvm_repro_syevj.cpp -o r && ./r
//
// This is a faithful, dependency-free extraction of rocSOLVER's small-matrix Jacobi
// eigensolver device function (run_syevj / syevj_small_kernel). It diagonalizes a batch
// of symmetric matrices; the accumulated eigenvector matrix A must come out orthogonal
// (A^T A = I). Build with -DBUG=1 (default) to use the upstream form, where the rotation
// (c,s) is produced by the tiy==0 lane, stored in LDS, and reloaded by every row-parallel
// lane after __syncthreads(). Build with -DBUG=0 to compute (c,s) per lane instead.
//
// On gfx1100 the -DBUG=1 build returns non-orthogonal eigenvectors (||A^T A - I|| = O(1),
// silently) once the 2D thread array exceeds ~512 threads (n >= ~40); -DBUG=0 is clean.
// The two forms are numerically identical (same (c,s), a pure function of unmodified Acpy
// entries), so the only difference is the post-barrier LDS reload -> a codegen defect.
#include
#include
#include
#include
#include
#ifndef BUG
#define BUG 1
#endif

static __device__ void givens(double f, double g, double& c, double& s, double& r){
if(g==0){ c=1; s=0; r=f; } else if(f==0){ c=0; s=1; r=g; }
else { r=std::hypot(f,g); double d=1.0/r; c=std::fabs(f)*d;
s=(f>=0? g*d : -g*d); if(f<0) r=-r; }
}

// Faithful transcription of run_syevj (evect=vector, uplo=upper, esort=none).
__global__ void __launch_bounds__(1024)
run_syevj_k(int n, double* AA, double* AcpyA, double* WW, int lda, double abstol,
int max_sweeps, int forced_ddy)
{
int tid=threadIdx.x, bid=blockIdx.z;
int even_n=n+n%2, half_n=even_n/2;
double* A = AA + (size_t)bid*lda*n;
double* Acpy= AcpyA + (size_t)bid*n*n;
double* W = WW + (size_t)bid*n;

// syevj_get_dims(n, 1024): ddy=min(256,half_n), ddx=min(1024/ddy,half_n); forced_ddy overrides ddy.
int ddy = forced_ddy>0 ? forced_ddy : (256i;j--){ aij=A[i+j*lda]; local_res+=2*aij*aij; Acpy[i+j*n]=aij; Acpy[j+i*n]=aij;
A[i+j*lda]=0; A[j+i*lda]=0; }
}
cosines_res[tix]=local_res; sines_diag[tix]=local_diag;
for(i=tix;itolerance){
for(int k=0;k0){ if(i==2||i==even_n-1) top[kx]=i-1; else top[kx]=i+((i%2==0)?-2:2); }
if(j==2||j==even_n-1) bottom[kx]=j-1; else bottom[kx]=j+((j%2==0)?-2:2);
}
__syncthreads();
}
}
if(tiy==0){ local_res=0; for(i=tix;i N(0,1);
// near-null SPD batch: A = Q diag(lam) Q^T, lam={1..1e-12, then 1e-16}
std::vector A((size_t)B*n*n), q(n*n), lam(n);
for(int k=0;k>>(n,dA,dAcpy,dW,n,1e-12,100,0);
hipDeviceSynchronize();
std::vector V(A.size()); hipMemcpy(V.data(),dA,sizeof(double)*A.size(),hipMemcpyDeviceToHost);
double worst=0; int nfail=0;
for(int b=0;bworst)worst=o; if(o>1e-8)nfail++; }
printf("n=%2d threads=%4d [%s] ||A^T A - I|| worst=%.3e corrupt=%d/%d %s\n",
n,threads, n<=58?"small":"blk", worst,nfail,B, nfail?" <== MISCOMPILE":" ok");
worst_fail+=nfail;
hipFree(dA);hipFree(dAcpy);hipFree(dW);
}
printf("\n%s\n", worst_fail? "REPRODUCED (non-orthogonal eigenvectors)":"clean");
return worst_fail?1:0;
}

```