#1929 [Fix] Build failure on gfx906 (Vega20 / Hygon DCU) due to missing STRATA_HIP_GFX906 guard in gr_upmix and hipblas types

本文

### Environment
- **Target Hardware**: Hygon DCU Z100 (amdgcn-amd-amdhsa--gfx906)
- **Toolchain**: DTK / HIP Clang 15
- **CMake Flags**: `-DSTRATA_HIP_GFX906=ON -DCMAKE_HIP_ARCHITECTURES=gfx906`

### Description
When building with `-DSTRATA_HIP_GFX906=ON`, compilation on `gfx906` targets hits three issues:

1. **Unconditional WMMA / __bf16 vector kernel in `src/prefill/kernels.cu`:**
   `gr_upmix_kernel` unconditionally defines and compiles `__bf16` vector types and `um_wmma` under `#if defined(__HIPCC__)`. Because `gfx906` does not feature hardware WMMA/MFMA or `__bf16` vector extensions, Clang fails with:
   ```text
   error: cannot compile this __bf16 vector builtin yet
   error: invalid instruction for target 'gfx906'
   ```
   **Suggested Fix:** Guard the WMMA kernel with `&& !defined(STRATA_HIP_GFX906)`, and provide a stub/no-op fallback returning `false` early in `gr_upmix()` for `STRATA_HIP_GFX906`.

2. **hipBLAS data type mismatch in `include/strata/platform/hip_compat/cublas_v2.h`:**
   `CUDA_R_16F` / `CUDA_R_32F` are defined as `HIP_R_16F` instead of `HIPBLAS_R_16F`. In certain HIP/ROCm/DTK releases, `hipblasGemmEx` strictly expects `hipblasDatatype_t`.

3. **Missing `cudaGraphUpload` fallback in `include/strata/platform/hip_compat/strata_hip.h`:**
   In older HIP runtimes / DTK (based on ROCm 5.7), `hipGraphUpload` is not available in `libamdhip64.so`. Adding an inline no-op fallback prevents link-time undefined reference errors.

### Diff
```diff
--- a/src/prefill/kernels.cu
+++ b/src/prefill/kernels.cu
@@ -1728,7 +1728,7 @@
-#if defined(__HIPCC__)
+#if defined(__HIPCC__) && !defined(STRATA_HIP_GFX906)
 typedef __bf16 um_b16 __attribute__((ext_vector_type(16)));
 typedef float um_f8 __attribute__((ext_vector_type(8)));
...
+#elif defined(STRATA_HIP_GFX906)
+__global__ void __launch_bounds__(256) gr_upmix_kernel(...) {}
+#endif
@@ -1851,7 +1857,11 @@
 bool gr_upmix(...) {
+#if defined(STRATA_HIP_GFX906)
+    return false;
+#elif defined(__HIPCC__)
```