Skip to content

Commit fa87b16

Browse files
committed
fix(turboquant): rebase hip-guard patch onto f2726891 (anchor moved)
The bump to turboquant f27268914 shifted the ggml-cuda.cu anchors the hip-guard patch targeted, so apply-patches failed on the copy2d hunk. Regenerated the patch against the new source: - Hunk 1 (ggml_cuda_copy2d_across_devices 3D-peer fast path) is now HIP-ported upstream in the fork itself (it wraps the peer path in #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) with a cudaMemcpy2DAsync #else fallback), so that hunk was DROPPED as now-upstream. - Hunk 2 (ggml_backend_cuda_device_event_new using plain cudaEventCreate) is still present upstream and still breaks the HIP shim, so it is kept and rebased onto the moved anchor (cudaEventCreateWithFlags + cudaEventDisableTiming). Verified with `git apply --check` and apply-patches.sh against a fresh turboquant @ f27268914 checkout: applies cleanly. Signed-off-by: Ettore Di Giacinto <mudler@localai.io> Assisted-by: Claude:opus-4.8 [Claude Code]
1 parent 1ba4633 commit fa87b16

1 file changed

Lines changed: 16 additions & 33 deletions

File tree

backend/cpp/turboquant/patches/0001-hip-guard-copy2d-peer-fastpath.patch

Lines changed: 16 additions & 33 deletions
Original file line numberDiff line numberDiff line change
@@ -5,51 +5,34 @@ that ggml's HIP (and MUSA) compatibility layer does not provide, breaking
55
the -gpu-rocm-hipblas-turboquant build:
66

77
1. ggml_cuda_copy2d_across_devices() (host-staged cross-device copy for
8-
split mul_mat output) uses the CUDA 3D-peer copy APIs
8+
split mul_mat output) used the CUDA 3D-peer copy APIs
99
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).
10+
cudaMemcpy3DPeerAsync, which HIP does not support. As of turboquant
11+
f27268914 the fork HIP-ports this itself: the peer fast path is now
12+
wrapped in #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) with
13+
an #else that does a plain cudaMemcpy2DAsync device-to-device copy. The
14+
hunk we used to carry here is therefore now upstream and has been
15+
DROPPED.
1716

18-
2. ggml_backend_cuda_device_event_new() creates its event with plain
17+
2. ggml_backend_cuda_device_event_new() still creates its event with plain
1918
cudaEventCreate, which ggml's HIP shim does not alias (it only aliases
20-
cudaEventCreateWithFlags). Use cudaEventCreateWithFlags(...,
19+
cudaEventCreateWithFlags). Use cudaEventCreateWithFlags(...,
2120
cudaEventDisableTiming) -- exactly what the rest of this file already
22-
does (cf. lines ~1034, ~3461) and HIP-safe.
21+
does (cf. lines ~1062, ~3523) and HIP-safe.
2322

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.
23+
CUDA builds are unaffected. Drop the remaining hunk once the fork HIP-ports
24+
it too; apply-patches.sh fails fast if an anchor goes stale.
2625

2726
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
28-
index 0427e6b..6352e6a 100644
27+
index 5d009f6d..622d28b9 100644
2928
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
3029
+++ 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_
30+
@@ -5795,7 +5795,7 @@ static ggml_backend_event_t ggml_backend_cuda_device_event_new(ggml_backend_dev_
4831
ggml_cuda_set_device(dev_ctx->device);
49-
32+
5033
cudaEvent_t event;
5134
- CUDA_CHECK(cudaEventCreate(&event));
5235
+ CUDA_CHECK(cudaEventCreateWithFlags(&event, cudaEventDisableTiming));
53-
36+
5437
return new ggml_backend_event {
5538
/* .device = */ dev,

0 commit comments

Comments
 (0)