Fix CI breaks from the off-graph KV stack (ROCm build, muse-glimmer offgraph e2e) (pytorch#23304)

Fixes CI breaks introduced by the off-graph KV stack (pytorch#23292–pytorch#23297).

### 1. ROCm build (`unittest-rocm-gfx950-rocm7.2`,
`test-voxtral-realtime-rocm-gfx950-rocm7.2`)
pytorch#23293 added `backends/cuda/runtime/cuda_kv_cache.h`, which includes
`<cuda_runtime.h>` directly. The ROCm build compiles `aoti_cuda_backend`
with `EXECUTORCH_USE_HIP` and no CUDA toolkit:

```
backends/cuda/runtime/cuda_kv_cache.h:14:10: fatal error: cuda_runtime.h: No such file or directory
```

- Include `executorch/extension/cuda/runtime_api.h` (the CUDA/HIP shim
used by the rest of the runtime) instead.
- Add HIP shims for the event APIs the KV cache uses: `cudaEvent_t`,
`cudaEventCreateWithFlags`, `cudaEventDisableTiming`,
`cudaEventDestroy`, `cudaEventRecord`, `cudaStreamWaitEvent`.

### 2. `test-muse-glimmer-cuda-e2e-kquant-17gb-solo-text-offgraph`
pytorch#23296 added this job and the `solo-text-offgraph` export branch, but
the early mode validation in `.ci/scripts/export_model_artifact.sh`
still only allowed `solo-text`/`dflash-image`:

```
Error: Muse Glimmer requires mode 'solo-text' or 'dflash-image'
```

Allow `solo-text-offgraph` there.

### 3. Off-graph export fails in AOTI lowering
Once past the mode check, the same job fails in export:

```
InductorError: RuntimeError: When autotune_at_compile_time is False, AOTInductor generates both JIT code and AOT code. The JIT code failed to run.
```

pytorch#23295 forced `autotune_at_compile_time=OFF` for the off-graph export.
With it off, AOTI tunes by running a JIT build of the wrapper on the
real constants. `LowerOffGraphKVPass` declares the KV storage constants
with zero-byte storage, because the runtime supplies the storage, so
nothing can run on them. The CUDA unit test only works because it swaps
them for real zero tensors first. This drops the override so AOTI keeps
its default compile-time autotune on generated inputs, which is the mode
the pass's `kv_len` derivation is written for.

Test plan: ROCm + CUDA CI on this PR.

Author: Gasoonjia

.ci/scripts/export_model_artifact.sh

@@ -289,8 +289,8 @@ if [ "$MODEL_NAME" = "muse_glimmer" ]; then
     echo "Error: Muse Glimmer requires quantization 'kquant-17gb' or 'kquant-dynamic'"
     exit 1
   fi
-  if [ "$MODE" != "solo-text" ] && [ "$MODE" != "dflash-image" ]; then
-    echo "Error: Muse Glimmer requires mode 'solo-text' or 'dflash-image'"
+  if [ "$MODE" != "solo-text" ] && [ "$MODE" != "solo-text-offgraph" ] && [ "$MODE" != "dflash-image" ]; then
+    echo "Error: Muse Glimmer requires mode 'solo-text', 'solo-text-offgraph' or 'dflash-image'"
     exit 1
   fi
 fi

backends/cuda/passes/lower_offgraph_kv.py

@@ -308,6 +308,8 @@ def __init__(self, manifest: dict[str, Any]) -> None:
     def _compile_storage(shape, dtype: torch.dtype, device):
         storage = torch.empty(shape, dtype=dtype, device=device)
         # AOTI needs the logical metadata, while the runtime supplies storage.
+        # Nothing can run on it, so lowering must keep AOTI's compile-time
+        # autotune (generated inputs) rather than the JIT pass on real constants.
         storage.untyped_storage().resize_(0)
         return storage
 

backends/cuda/runtime/cuda_kv_cache.h

@@ -13,10 +13,9 @@
 #include <optional>
 #include <string_view>
 
-#include <cuda_runtime.h>
-
 #include <executorch/backends/cuda/runtime/backend_options.h>
 #include <executorch/backends/cuda/runtime/cuda_delegate_handle.h>
+#include <executorch/extension/cuda/runtime_api.h>
 #include <executorch/extension/llm/cache/cache.h>
 #include <executorch/runtime/core/error.h>
 #include <executorch/runtime/core/evalue.h>

examples/models/muse-glimmer/export/export_solo.py

@@ -316,7 +316,6 @@ def _partitioner_for(name: str) -> "CudaPartitioner":
                     # the runtime cannot recover the width from the graph; every
                     # method here takes input_pos second, shaped [T].
                     CompileSpec(OFFGRAPH_KV_STEP_WIDTH_COMPILE_SPEC, b"1:0"),
-                    CompileSpec("autotune_at_compile_time", b"OFF"),
                 )
             )
         return CudaPartitioner(compile_specs)

extension/cuda/runtime_api.h

@@ -15,6 +15,7 @@
 #include <hip/hip_runtime.h>
 
 using cudaError_t = hipError_t;
+using cudaEvent_t = hipEvent_t;
 using cudaGraph_t = hipGraph_t;
 using cudaGraphExec_t = hipGraphExec_t;
 using cudaMemcpyKind = hipMemcpyKind;
@@ -23,6 +24,7 @@ using cudaStreamCaptureMode = hipStreamCaptureMode;
 using cudaStream_t = hipStream_t;
 
 inline constexpr cudaError_t cudaSuccess = hipSuccess;
+inline constexpr unsigned int cudaEventDisableTiming = hipEventDisableTiming;
 inline constexpr cudaMemcpyKind cudaMemcpyDeviceToDevice =
     hipMemcpyDeviceToDevice;
 inline constexpr cudaMemcpyKind cudaMemcpyDeviceToHost = hipMemcpyDeviceToHost;
@@ -50,6 +52,20 @@ inline cudaError_t cudaDeviceSynchronize() {
   return hipDeviceSynchronize();
 }
 
+inline cudaError_t cudaEventCreateWithFlags(
+    cudaEvent_t* event,
+    unsigned int flags) {
+  return hipEventCreateWithFlags(event, flags);
+}
+
+inline cudaError_t cudaEventDestroy(cudaEvent_t event) {
+  return hipEventDestroy(event);
+}
+
+inline cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream) {
+  return hipEventRecord(event, stream);
+}
+
 inline cudaError_t cudaFree(void* ptr) {
   return hipFree(ptr);
 }
@@ -167,6 +183,13 @@ inline cudaError_t cudaStreamSynchronize(cudaStream_t stream) {
   return hipStreamSynchronize(stream);
 }
 
+inline cudaError_t cudaStreamWaitEvent(
+    cudaStream_t stream,
+    cudaEvent_t event,
+    unsigned int flags) {
+  return hipStreamWaitEvent(stream, event, flags);
+}
+
 #else
 
 #include <cuda_runtime.h>