[AMDGPU] misses SADDR+VADDR load for a loop-carried uniform base pointer (gfx1151)

## Summary

amdclang++/hipcc fails to select the `global_load ... SADDR+VADDR` form when it
is applicable.
Instead it keeps the full 64-bit address as a
per-lane VGPR induction variable and emits a pure-VGPR `global_load ... off`,
advancing the address with VALU ops instead of SALU ops per iteration.
Using VALU operations for index updates occupies the VALU unit and thus steals
time from VALU ops like `v_wmma` or `v_dot` for peak GEMM performance.

Tested with the ROCm 7.14 amdclang++ (`--offload-arch=gfx1151 -O3`).

## Reproducer

```c++
__global__ void saddr_loop(const __half* __restrict__ in,
float* __restrict__ out, int K) {
const __half* base = in; // uniform
float acc = 0.f;
for (int i = 0; i < K; ++i) {
acc += __half2float(base[threadIdx.x]); // base (uniform) + threadIdx.x
base += 4; // loop-carried uniform advance
}
out[threadIdx.x] = acc;
}
```

ISA (gfx1151, -O3): pure-VGPR address, base advanced in VGPRs:

```asm
global_load_d16_b16 v4, v[1:2], off ; SADDR = null
v_add_co_u32 v1, vcc_lo, v1, 8 ; advance 64-bit base per iter
v_add_co_ci_u32_e64 v2, null, 0, v2, vcc_lo
```

## Alternative reproducer

```c++
__global__ void saddr_canon(const __half* __restrict__ in,
float* __restrict__ out, int K) {
float acc = 0.f;
for (int i = 0; i < K; ++i)
acc += __half2float(in[threadIdx.x + 32*i]);
out[threadIdx.x] = acc;
}
```

ISA: same pure-VGPR loop, address advanced by 64 bytes per iteration:

```asm
global_load_d16_b16 v5, v[2:3], off
v_add_co_u32 v2, vcc_lo, v2, 64
v_add_co_ci_u32_e64 v3, null, 0, v3, vcc_lo
```

## Expected
Keep `base` in an SGPR pair, advance with a scalar add, and emit
`global_load_d16_b16 vdst, v_tid_off, s[base:base+1]`.

```asm
global_load_d16_b16 v4, v_tid_off, s[0:1] ; SADDR = s[0:1]
s_add_u32 s0, s0, 8 ; scalar advance lo
s_addc_u32 s1, s1, 0 ; hi
```

Start by compiling both supplied C++ reproducers with amdclang++/hipcc for gfx1151 at -O3 and compare the generated ISA with the expected SGPR-based form. Trace the AMDGPU compiler code path handling loop-carried uniform pointers and address selection; done means the base remains in an SGPR pair, advances with scalar adds, and the load uses SADDR plus the per-lane offset.

Written by the indexing model from the issue text.