From cebaa4da00d13e6f07edc2e25d2bcbab2e4da55c Mon Sep 17 00:00:00 2001 From: Davide Cifarelli Date: Sun, 7 Jun 2026 16:23:52 +0200 Subject: [PATCH] fix(ggml-cuda): make fattn-sparse build under HIP fattn-sparse.cu used cudaStreamDefault, which is not defined in the HIP translation (the HIP shim does not alias it), so the ggml-hip build failed with "use of undeclared identifier cudaStreamDefault". Use ((cudaStream_t)0) (the default-stream handle), valid and identical under both CUDA and HIP. Co-Authored-By: WOZCODE --- ggml/src/ggml-cuda/fattn-sparse.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/fattn-sparse.cu b/ggml/src/ggml-cuda/fattn-sparse.cu index 728e03d9caf2..151578542f20 100644 --- a/ggml/src/ggml-cuda/fattn-sparse.cu +++ b/ggml/src/ggml-cuda/fattn-sparse.cu @@ -206,14 +206,14 @@ void ggml_cuda_flash_attn_sparse(ggml_backend_cuda_context & ctx, ggml_tensor * cudaEvent_t ev_conv; CUDA_CHECK(cudaEventCreateWithFlags(&ev_conv, cudaEventDisableTiming)); CUDA_CHECK(cudaEventRecord(ev_conv, stream)); - CUDA_CHECK(cudaStreamWaitEvent(cudaStreamDefault, ev_conv, 0)); + CUDA_CHECK(cudaStreamWaitEvent(((cudaStream_t)0), ev_conv, 0)); CUDA_CHECK(cudaEventDestroy(ev_conv)); int err = s_sparse_kernel(Q_pf, K_pf, V_pf, O_pf, B, S, H, Hk, D, scale, alpha); GGML_ASSERT(err == 0 && "sparse attention kernel failed"); cudaEvent_t ev_pf; CUDA_CHECK(cudaEventCreateWithFlags(&ev_pf, cudaEventDisableTiming)); - CUDA_CHECK(cudaEventRecord(ev_pf, cudaStreamDefault)); + CUDA_CHECK(cudaEventRecord(ev_pf, ((cudaStream_t)0))); CUDA_CHECK(cudaStreamWaitEvent(stream, ev_pf, 0)); CUDA_CHECK(cudaEventDestroy(ev_pf));