[AMDGPU] gfx1250 WMMA co-execution WAR hazard is not applied to the C (src2) operand, miscompiling v_wmma_bf16f32_16x16x32_bf16

### Summary

On gfx1250 the WMMA co-execution hazard handling added in #149865 covers the WAR case "WMMA reads (A0/B0/Index0), VALU writes (D1)" but never considers the WMMA's **C matrix** (`src2` of a non-SWMMAC WMMA). A VALU instruction scheduled into the MMA's shadow can therefore overwrite a register the in-flight WMMA is still reading as C, and the MMA consumes the new value.

This is latent for almost every WMMA, because C is normally tied to `vdst` (in-place accumulate) and so is covered incidentally by the existing `D0` check. It becomes a miscompile for the mixed-C/D instructions, whose D is a different, narrower register tuple than C. `v_wmma_bf16f32_16x16x32_bf16` is the only such instruction in the ISA today (`int_amdgcn_wmma_bf16f32_16x16x32_bf16` is the sole user of `AMDGPUWmmaIntrinsicModsCDiff` in `IntrinsicsAMDGPU.td`), and it produces wrong results.

### Environment

- Target: `gfx1250`, wave32
- Reproduced with AMD clang 23.0.0git (ROCm/llvm-project `2a67d831a09cc538194bd595279ee91258570c10`, 2026-07-27)
- Present at `-O2` and `-O3`; not at `-O0`/`-O1` (the scheduler does not fill the MMA's shadow there)
- Still present in `llvm/llvm-project` `main` at `18b9d7282c067d83edda8006c25542d9490e3c6e` (LLVM 24.0.0git, 2026-09-11) — see Reproducer 1

### Reproducer 1: the hazard recognizer in isolation

This is the one to look at first. It needs no GPU, and unlike the two below it does not depend on what the register allocator happens to do, so it is the version that still demonstrates the omission on `main` today. The two functions hold the same `V_WMMA_BF16F32_16X16X32_BF16_w32_twoaddr` followed by one VALU write at the same distance; only the register being written differs. In the first it is `$vgpr8`, which belongs to C (`$vgpr2_..._$vgpr9`); in the second it is `$vgpr10`, which belongs to A/B. The names follow the existing cases in `llvm/test/CodeGen/AMDGPU/wmma-coexecution-valu-hazards.mir`, and the WMMA line is verbatim from `-stop-after=postrapseudos` on Reproducer 2.

```mir
# llc -mtriple=amdgpu12.50 -run-pass post-RA-hazard-rec c_operand_hazard.mir -o -
---
name: test_wmma_bf16f32_16x16x32_bf16_C0_overlaps_D1
body: |
bb.0:
early-clobber $vgpr18_vgpr19_vgpr20_vgpr21 = V_WMMA_BF16F32_16X16X32_BF16_w32_twoaddr killed $vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17, $vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17, 8, killed $vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9, 0, 0, 0, 0, implicit $exec
$vgpr8 = V_LSHLREV_B32_e32 5, $vgpr0, implicit $exec
...
---
name: test_wmma_bf16f32_16x16x32_bf16_A0_overlaps_D1
body: |
bb.0:
early-clobber $vgpr18_vgpr19_vgpr20_vgpr21 = V_WMMA_BF16F32_16X16X32_BF16_w32_twoaddr killed $vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17, $vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17, 8, killed $vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9, 0, 0, 0, 0, implicit $exec
$vgpr10 = V_LSHLREV_B32_e32 5, $vgpr0, implicit $exec
...
```

The `A0` case comes back with four `V_NOP_e32` inserted before the `V_LSHLREV_B32_e32`. The `C0` case comes back with none.

### Reproducer 2: LLVM IR through llc

Read the caveat first: **this IR no longer miscompiles on `main`**, so please do not take it as the check for whether the bug is present — Reproducer 1 is. It is included because it is what the miscompile actually looked like in the wild, and because what `main` does with it is itself part of the argument.

On `main` at `18b9d7282c` the allocator puts C in `v[0:7]` and D in `v[16:19]`, and the result is safe purely by accident:

```asm
v_wmma_bf16f32_16x16x32_bf16 v[16:19], v[8:15], v[8:15], v[0:7]
v_nop
v_nop
v_nop
v_nop
v_and_b32_e32 v1, 0xffff0000, v16 ; writes C's dword 1
v_lshlrev_b32_e32 v0, 16, v16 ; writes C's dword 0
```

Those four `V_NOP`s were inserted for the `D0` read hazard on `v16`, which has nothing to do with C. They happen to supply exactly the four VALU wait states that the C write also needs, so the margin here is zero: any scheduling or allocation change that drops one of them, or that moves the first write to C one slot earlier, turns this back into a miscompile.

The IR:

```llvm
; llc -mtriple=amdgpu12.50-amd-amdhsa -O3 repro.ll -o -

target triple = "amdgcn-amd-amdhsa"

declare <8 x bfloat> @llvm.amdgcn.wmma.bf16f32.16x16x32.bf16.v8bf16.v16bf16.v8f32(<16 x bfloat>, <16 x bfloat>, i16 immarg, <8 x float>, i1 immarg, i1 immarg)
declare i32 @llvm.amdgcn.workitem.id.x()

define amdgpu_kernel void @repro(ptr addrspace(1) writeonly captures(none) %out) #0 {
%tid = tail call i32 @llvm.amdgcn.workitem.id.x()
%base = shl nuw nsw i32 %tid, 3

; C[j] = lane * 8 + j, for j in [0, 8)
%i0 = insertelement <8 x i32> poison, i32 %base, i64 0
%spl = shufflevector <8 x i32> %i0, <8 x i32> poison, <8 x i32> zeroinitializer
%idx = or disjoint <8 x i32> %spl,
%c = uitofp nneg <8 x i32> %idx to <8 x float>

; A = B = 0, so D must come back equal to C, rounded to bf16.
%d = tail call <8 x bfloat> @llvm.amdgcn.wmma.bf16f32.16x16x32.bf16.v8bf16.v16bf16.v8f32(
<16 x bfloat> zeroinitializer, <16 x bfloat> zeroinitializer, i16 0,
<8 x float> %c, i1 false, i1 false)

; Store D back as 4 pairs of floats. The address arithmetic for these stores is what the
; scheduler places in the MMA's shadow.
%o0 = zext nneg i32 %base to i64
%p0 = getelementptr inbounds nuw [4 x i8], ptr addrspace(1) %out, i64 %o0
%d01 = shufflevector <8 x bfloat> %d, <8 x bfloat> poison, <2 x i32>
%f01 = fpext <2 x bfloat> %d01 to <2 x float>
store <2 x float> %f01, ptr addrspace(1) %p0, align 4

%b2 = or disjoint i32 %base, 2
%o2 = zext nneg i32 %b2 to i64
%p2 = getelementptr inbounds nuw [4 x i8], ptr addrspace(1) %out, i64 %o2
%d23 = shufflevector <8 x bfloat> %d, <8 x bfloat> poison, <2 x i32>
%f23 = fpext <2 x bfloat> %d23 to <2 x float>
store <2 x float> %f23, ptr addrspace(1) %p2, align 4

%b4 = or disjoint i32 %base, 4
%o4 = zext nneg i32 %b4 to i64
%p4 = getelementptr inbounds nuw [4 x i8], ptr addrspace(1) %out, i64 %o4
%d45 = shufflevector <8 x bfloat> %d, <8 x bfloat> poison, <2 x i32>
%f45 = fpext <2 x bfloat> %d45 to <2 x float>
store <2 x float> %f45, ptr addrspace(1) %p4, align 4

%b6 = or disjoint i32 %base, 6
%o6 = zext nneg i32 %b6 to i64
%p6 = getelementptr inbounds nuw [4 x i8], ptr addrspace(1) %out, i64 %o6
%d67 = shufflevector <8 x bfloat> %d, <8 x bfloat> poison, <2 x i32>
%f67 = fpext <2 x bfloat> %d67 to <2 x float>
store <2 x float> %f67, ptr addrspace(1) %p6, align 4

ret void
}

attributes #0 = { nounwind "target-cpu"="gfx1250" }
```

The ISA this produced with AMD clang 23.0.0git, where the allocator chose C = `v[2:9]` and D = `v[18:21]` (trimmed; `-O3`):

```asm
v_cvt_f32_u32_e32 v9, v6 ; C[7]
v_cvt_f32_u32_e32 v8, v7 ; C[6]
...
v_wmma_bf16f32_16x16x32_bf16 v[18:21], v[10:17], v[10:17], v[2:9]
v_lshlrev_b32_e32 v8, 5, v0 ; <-- overwrites v8, i.e. C's dword 6
v_nop
v_nop
v_nop
v_and_b32_e32 v1, 0xffff0000, v18
v_lshlrev_b32_e32 v0, 16, v18
```

C is `v[2:9]`, so `v8` holds accumulator element 6. The `v_lshlrev_b32_e32 v8, 5, v0` (store-address arithmetic, unrelated to the MMA) is issued in the very first VALU slot after the MMA, while the hardware is still reading C. The three `V_NOP`s that the hazard recognizer did insert are for the `D0`-read hazard and are placed *after* the clobber.

Note the asymmetry: `v9` (C's element 7) is also reused for address arithmetic, but only after those three `V_NOP`s, and it survives. Which element is corrupted tracks exactly how many VALU slots separate the MMA from the write.

### Reproducer 3: runtime check on hardware

Built with AMD clang 23.0.0git, i.e. the allocation shown in Reproducer 2. This is where the wrong numbers were actually observed; it exits non-zero when it reproduces.

```c++
#include
#include

typedef __bf16 bf16x16 __attribute__((ext_vector_type(16)));
typedef __bf16 bf16x8 __attribute__((ext_vector_type(8)));
typedef float floatx8 __attribute__((ext_vector_type(8)));

// A = B = 0, so D must come back equal to C.
__global__ void k(float* out)
{
const int lane = threadIdx.x;
bf16x16 a = {}, b = {};
floatx8 c;
for(int j = 0; j < 8; j++)
c[j] = float(lane * 8 + j);

bf16x8 d = __builtin_amdgcn_wmma_bf16f32_16x16x32_bf16(false, a, false, b, 0, c, false, false);
for(int j = 0; j < 8; j++)
out[lane * 8 + j] = float(d[j]);
}

int main()
{
float* d_out;
(void)hipMalloc(&d_out, 32 * 8 * sizeof(float));
(void)hipMemset(d_out, 0, 32 * 8 * sizeof(float));
k<<<1, 32>>>(d_out);
(void)hipDeviceSynchronize();

float h[32 * 8];
(void)hipMemcpy(h, d_out, sizeof(h), hipMemcpyDeviceToHost);

int bad = 0;
for(int i = 0; i < 32 * 8; i++)
bad += (h[i] != float(i));
printf("lane 0 D:");
for(int j = 0; j < 8; j++)
printf(" %4.1f", h[j]);
printf(" (expected 0..7, %d of 256 elements wrong)\n", bad);
return bad != 0;
}
```

```
$ hipcc --offload-arch=gfx1250 -O3 minimal.hip -o minimal && ./minimal
lane 0 D: 0.0 1.0 2.0 3.0 4.0 5.0 0.0 7.0 (expected 0..7, 32 of 256 elements wrong)
```

Element 6 reads 0 on all 32 lanes; every other element is correct. The clobbering value is `lane * 32`, which as an f32 bit pattern is a denormal that rounds to 0 in bf16, and is exactly 0 for lane 0.

### Evidence that it is the C read, not the D write

Three variations pin it down:

1. Copying `D` out with `__builtin_memcpy` instead of converting it element-wise gives a different register allocation in which nothing clobbers C, and all four D dwords are then correct (`3f800000 40404000 40a04080 40e040c0`, i.e. `0.0 1.0 ... 6.0 7.0`). So the instruction and the hardware are fine.
2. `__builtin_amdgcn_sched_barrier(0)` after the intrinsic makes the compiler emit four `v_nop`s before reusing `v8`, and the result is correct on all 8 elements. Four matches `VALUWaitStates[0] = 4` for this instruction's category.
3. Replacing the barrier with `asm volatile("s_nop 7\n\ts_nop 7")` does *not* help, and makes **both** elements 6 and 7 wrong: the scheduler puts the `v8` and `v9` writes ahead of the `s_nop`s, and `s_nop` is not a VALU slot so it does not satisfy the requirement either.

Start with llvm/test/CodeGen/AMDGPU/wmma-coexecution-valu-hazards.mir and run the post-RA-hazard-rec check using the C0 and A0 reproducer cases. Trace the existing WMMA co-execution hazard handling and verify that the C operand case receives the required VALU wait states, with the test showing the expected V_NOP instructions.

Written by the indexing model from the issue text.