Skip to content

Commit bd8a735

Browse files
committed
fix(turboquant): re-anchor HIP event patch for fork a33ef00b (mudler#10235)
The fork a33ef00b HIP-ported ggml_cuda_copy2d_across_devices() itself (it now guards the cudaMemcpy3DPeer fast path with #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) and falls back to a plain 2D device-to-device copy), so the former copy2d hunk's anchors went stale. Retire that hunk and re-anchor the still-needed event hunk: ggml_backend_cuda_device_event_new() still uses plain cudaEventCreate, which ggml's HIP shim does not alias, so keep the cudaEventCreateWithFlags (cudaEventDisableTiming) rewrite. Validated with apply-patches.sh against a fresh clone at the pinned commit. Signed-off-by: Ettore Di Giacinto <mudler@localai.io> Assisted-by: Claude:claude-opus-4-8 [Claude Code]
1 parent 787c635 commit bd8a735

1 file changed

Lines changed: 18 additions & 40 deletions

File tree

Lines changed: 18 additions & 40 deletions
Original file line numberDiff line numberDiff line change
@@ -1,55 +1,33 @@
11
hip: port the turboquant CUDA additions that ggml's HIP shim doesn't cover
22

3-
The turboquant fork adds/modifies a few ggml-cuda.cu spots with CUDA APIs
4-
that ggml's HIP (and MUSA) compatibility layer does not provide, breaking
5-
the -gpu-rocm-hipblas-turboquant build:
3+
The turboquant fork adds/modifies ggml-cuda.cu spots with CUDA APIs that
4+
ggml's HIP (and MUSA) compatibility layer does not provide, breaking the
5+
-gpu-rocm-hipblas-turboquant build:
66

7-
1. ggml_cuda_copy2d_across_devices() (host-staged cross-device copy for
8-
split mul_mat output) uses the CUDA 3D-peer copy APIs
9-
cudaMemcpy3DPeerParms / make_cudaPitchedPtr / make_cudaExtent /
10-
cudaMemcpy3DPeerAsync. HIP genuinely does not support these (see the
11-
fork's own comment "HIP does not support cudaMemcpy3DPeerAsync"), so
12-
guard the peer fast path with #if !defined(GGML_USE_HIP) &&
13-
!defined(GGML_USE_MUSA) -- matching how the fork already guards the
14-
same API for the sibling 2D copy -- and fall through to the existing
15-
cudaMemcpyAsync staging fallback below (functionally identical,
16-
slightly slower on multi-GPU ROCm).
7+
ggml_backend_cuda_device_event_new() creates its event with plain
8+
cudaEventCreate, which ggml's HIP shim does not alias (it only aliases
9+
cudaEventCreateWithFlags). Use cudaEventCreateWithFlags(...,
10+
cudaEventDisableTiming) -- exactly what the rest of this file already
11+
does (cf. ggml_cuda_set_main_device extras, copy_event) and HIP-safe.
1712

18-
2. ggml_backend_cuda_device_event_new() creates its event with plain
19-
cudaEventCreate, which ggml's HIP shim does not alias (it only aliases
20-
cudaEventCreateWithFlags). Use cudaEventCreateWithFlags(...,
21-
cudaEventDisableTiming) -- exactly what the rest of this file already
22-
does (cf. lines ~1034, ~3461) and HIP-safe.
13+
CUDA builds are unaffected. Drop this patch once the fork HIP-ports the
14+
event create; apply-patches.sh fails fast if the anchor goes stale.
2315

24-
CUDA builds are unaffected. Drop the relevant hunk once the fork HIP-ports
25-
these; apply-patches.sh fails fast if an anchor goes stale.
16+
Note: the fork has since HIP-ported ggml_cuda_copy2d_across_devices()
17+
itself (it now guards the cudaMemcpy3DPeer fast path with
18+
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) and falls back to a
19+
plain 2D device-to-device copy), so the former copy2d hunk was retired.
2620

2721
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
28-
index 0427e6b..6352e6a 100644
22+
index 08cfd37f..b97f0012 100644
2923
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
3024
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
31-
@@ -1933,6 +1933,7 @@ static cudaError_t ggml_cuda_copy2d_across_devices(
32-
size_t width, size_t height, cudaStream_t dst_stream, cudaStream_t src_stream) {
33-
34-
const auto & info = ggml_cuda_info();
35-
+#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) // 3D-peer copy types unmapped by ggml's HIP/MUSA shim; use staging fallback below
36-
if (info.peer_access[src_device][dst_device]) {
37-
cudaMemcpy3DPeerParms p = {};
38-
p.dstDevice = dst_device;
39-
@@ -1942,6 +1943,7 @@ static cudaError_t ggml_cuda_copy2d_across_devices(
40-
p.extent = make_cudaExtent(width, height, 1);
41-
return cudaMemcpy3DPeerAsync(&p, dst_stream);
42-
}
43-
+#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
44-
45-
// Fallback: stage all rows through a single contiguous pinned buffer
46-
int prev_device = ggml_cuda_get_device();
47-
@@ -5714,7 +5716,7 @@ static ggml_backend_event_t ggml_backend_cuda_device_event_new(ggml_backend_dev_
25+
@@ -5794,7 +5794,7 @@ static ggml_backend_event_t ggml_backend_cuda_device_event_new(ggml_backend_dev_
4826
ggml_cuda_set_device(dev_ctx->device);
49-
27+
5028
cudaEvent_t event;
5129
- CUDA_CHECK(cudaEventCreate(&event));
5230
+ CUDA_CHECK(cudaEventCreateWithFlags(&event, cudaEventDisableTiming));
53-
31+
5432
return new ggml_backend_event {
5533
/* .device = */ dev,

0 commit comments

Comments
 (0)