|
1 | 1 | hip: port the turboquant CUDA additions that ggml's HIP shim doesn't cover |
2 | 2 |
|
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 creates backend events with plain cudaEventCreate, |
| 4 | +which ggml's HIP shim does not alias (it only aliases |
| 5 | +cudaEventCreateWithFlags). Use cudaEventCreateWithFlags(..., |
| 6 | +cudaEventDisableTiming), exactly as the rest of this file does. |
6 | 7 |
|
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). |
17 | | - |
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. |
23 | | - |
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. |
| 8 | +CUDA builds are unaffected. Drop this patch once the fork HIP-ports the |
| 9 | +event creation; apply-patches.sh fails fast if the anchor goes stale. |
26 | 10 |
|
27 | 11 | diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu |
28 | | -index 0427e6b..6352e6a 100644 |
| 12 | +index 7d35c1a..2908acb 100644 |
29 | 13 | --- a/ggml/src/ggml-cuda/ggml-cuda.cu |
30 | 14 | +++ 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_ |
| 15 | +@@ -5795,7 +5795,7 @@ static ggml_backend_event_t ggml_backend_cuda_device_event_new(ggml_backend_dev_ |
48 | 16 | ggml_cuda_set_device(dev_ctx->device); |
49 | 17 |
|
50 | 18 | cudaEvent_t event; |
|
0 commit comments