Skip to content

Commit 7402d1f

Browse files
localai-botmudler
andauthored
chore(turboquant): bump to 7d9715f1 + fix compilation against rebased fork (#10205)
* chore(turboquant): bump TheTom/llama-cpp-turboquant to 7d9715f1 Signed-off-by: Ettore Di Giacinto <mudler@localai.io> Assisted-by: Claude:claude-opus-4-8 [Claude Code] * fix(turboquant): drop obsolete legacy-spec shim after fork rebased The TheTom/llama-cpp-turboquant fork (pin c9aa86a) rebased past the upstream common_params_speculative refactor (ggml-org/llama.cpp #22397/#22838/#22964), the model_tgt rename (#22838) and get_media_marker (#21962). The old fork-compat shim forced now-wrong legacy code paths, breaking the build with errors like 'struct common_params_speculative has no member named mparams_dft / type' and 'server_context_impl has no member named model'. Remove the obsolete LOCALAI_LEGACY_LLAMA_CPP_SPEC branches from the shared grpc-server.cpp (stock llama-cpp and the modern fork both take the modern path now), and narrow the one remaining gap (the fork still lacks common_params::checkpoint_min_step) to a dedicated LOCALAI_TURBOQUANT_NO_CHECKPOINT_MIN_STEP guard injected by patch-grpc-server.sh. The patch script now only adds the turbo2/3/4 KV-cache types and injects that one macro. Signed-off-by: Ettore Di Giacinto <mudler@localai.io> Assisted-by: Claude:claude-opus-4-8 [Claude Code] * fix(turboquant): HIP-port the fork's CUDA additions (copy2d 3D-peer + cudaEventCreate) The turboquant fork adds/modifies a few ggml-cuda.cu spots with CUDA APIs that ggml's HIP/MUSA shim does not provide, breaking the -gpu-rocm-hipblas-turboquant build. patches/0001-hip-guard-copy2d-peer-fastpath.patch (applied by apply-patches.sh) ports them: - Guard ggml_cuda_copy2d_across_devices's 3D-peer copy fast path with #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) so HIP/MUSA fall through to the existing cudaMemcpyAsync staging fallback (HIP genuinely lacks cudaMemcpy3DPeerAsync, per the fork's own comment). - Create the device event in ggml_backend_cuda_device_event_new with the HIP-aliased cudaEventCreateWithFlags(.., cudaEventDisableTiming) instead of the un-aliased plain cudaEventCreate, matching this file's own usage elsewhere. CUDA builds are unaffected. Signed-off-by: Ettore Di Giacinto <mudler@localai.io> Assisted-by: Claude:claude-opus-4-8 [Claude Code] * ci(turboquant): drop the ROCm/hipblas build flavor The TheTom/llama-cpp-turboquant fork is not ROCm-clean at the current pin: beyond the CUDA-API gaps already patched (3D-peer copy, cudaEventCreate), its llama.cpp base fails to compile the flash-attention MMA f16 kernels for head-dim 640 under HIP (cols_per_warp evaluates to 0 -> division-by-zero / non-constant static asserts in fattn-mma-f16.cuh). That is a deep ggml-on-ROCm kernel issue, not something a small fork patch can paper over. Drop -gpu-rocm-hipblas-turboquant from the build matrix so turboquant still ships for cpu / cublas / vulkan / sycl. Re-add it once the fork's HIP path compiles (or upstream ggml fixes the large-head-dim MMA kernels for ROCm). Signed-off-by: Ettore Di Giacinto <mudler@localai.io> Assisted-by: Claude:claude-opus-4-8 [Claude Code] --------- Signed-off-by: Ettore Di Giacinto <mudler@localai.io> Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
1 parent 8c42695 commit 7402d1f

5 files changed

Lines changed: 91 additions & 145 deletions

File tree

.github/backend-matrix.yml

Lines changed: 0 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -1766,20 +1766,6 @@ include:
17661766
dockerfile: "./backend/Dockerfile.llama-cpp"
17671767
context: "./"
17681768
ubuntu-version: '2404'
1769-
- build-type: 'hipblas'
1770-
cuda-major-version: ""
1771-
cuda-minor-version: ""
1772-
platforms: 'linux/amd64'
1773-
tag-latest: 'auto'
1774-
tag-suffix: '-gpu-rocm-hipblas-turboquant'
1775-
builder-base-image: 'quay.io/go-skynet/ci-cache:base-grpc-rocm-amd64'
1776-
runs-on: 'ubuntu-latest'
1777-
base-image: "rocm/dev-ubuntu-24.04:7.2.1"
1778-
skip-drivers: 'false'
1779-
backend: "turboquant"
1780-
dockerfile: "./backend/Dockerfile.turboquant"
1781-
context: "./"
1782-
ubuntu-version: '2404'
17831769
- build-type: 'hipblas'
17841770
cuda-major-version: ""
17851771
cuda-minor-version: ""

backend/cpp/llama-cpp/grpc-server.cpp

Lines changed: 9 additions & 50 deletions
Original file line numberDiff line numberDiff line change
@@ -482,23 +482,13 @@ static void params_parse(server_context& /*ctx_server*/, const backend::ModelOpt
482482
if (!request->draftmodel().empty()) {
483483
params.speculative.draft.mparams.path = request->draftmodel();
484484
// Default to draft type if a draft model is set but no explicit type.
485-
// Upstream (post ggml-org/llama.cpp#22838) made the speculative type a
486-
// vector; the turboquant fork still uses the legacy scalar. The
487-
// LOCALAI_LEGACY_LLAMA_CPP_SPEC macro is injected by
488-
// backend/cpp/turboquant/patch-grpc-server.sh for fork builds only.
489-
// Upstream renamed COMMON_SPECULATIVE_TYPE_DRAFT -> ..._DRAFT_SIMPLE
490-
// in ggml-org/llama.cpp#22964; the fork still uses the old name.
491-
#ifdef LOCALAI_LEGACY_LLAMA_CPP_SPEC
492-
if (params.speculative.type == COMMON_SPECULATIVE_TYPE_NONE) {
493-
params.speculative.type = COMMON_SPECULATIVE_TYPE_DRAFT;
494-
}
495-
#else
485+
// Upstream made the speculative type a vector (ggml-org/llama.cpp#22838)
486+
// and renamed COMMON_SPECULATIVE_TYPE_DRAFT -> ..._DRAFT_SIMPLE (#22964).
496487
const bool no_spec_type = params.speculative.types.empty() ||
497488
(params.speculative.types.size() == 1 && params.speculative.types[0] == COMMON_SPECULATIVE_TYPE_NONE);
498489
if (no_spec_type) {
499490
params.speculative.types = { COMMON_SPECULATIVE_TYPE_DRAFT_SIMPLE };
500491
}
501-
#endif
502492
}
503493

504494
// params.model_alias ??
@@ -574,9 +564,10 @@ static void params_parse(server_context& /*ctx_server*/, const backend::ModelOpt
574564
// tokens (0 disables the minimum). Match upstream's default (256). This
575565
// field was renamed from `checkpoint_every_nt` in llama.cpp; the semantics
576566
// also shifted from a fixed cadence to a minimum spacing. The turboquant
577-
// fork branched before the field existed, so skip it on the legacy path
578-
// (LOCALAI_LEGACY_LLAMA_CPP_SPEC is injected by patch-grpc-server.sh).
579-
#ifndef LOCALAI_LEGACY_LLAMA_CPP_SPEC
567+
// fork still lacks common_params::checkpoint_min_step, so skip it there
568+
// (LOCALAI_TURBOQUANT_NO_CHECKPOINT_MIN_STEP is injected by
569+
// backend/cpp/turboquant/patch-grpc-server.sh).
570+
#ifndef LOCALAI_TURBOQUANT_NO_CHECKPOINT_MIN_STEP
580571
params.checkpoint_min_step = 256;
581572
#endif
582573

@@ -752,7 +743,7 @@ static void params_parse(server_context& /*ctx_server*/, const backend::ModelOpt
752743
params.cache_idle_slots = false;
753744
}
754745

755-
#ifndef LOCALAI_LEGACY_LLAMA_CPP_SPEC
746+
#ifndef LOCALAI_TURBOQUANT_NO_CHECKPOINT_MIN_STEP
756747
// --- minimum context-checkpoint spacing (upstream -cms / --checkpoint-min-step) ---
757748
// 0 disables the minimum-spacing gate. Old option names (`checkpoint_every_nt`,
758749
// `checkpoint_every_n_tokens`) are kept as aliases for backward compatibility
@@ -906,17 +897,6 @@ static void params_parse(server_context& /*ctx_server*/, const backend::ModelOpt
906897

907898
// Speculative decoding options
908899
} else if (!strcmp(optname, "spec_type") || !strcmp(optname, "speculative_type")) {
909-
#ifdef LOCALAI_LEGACY_LLAMA_CPP_SPEC
910-
// Fork only knows a single scalar `type`. Take the first comma-
911-
// separated value and assign it via the singular helper.
912-
std::string first = optval_str;
913-
const auto comma = first.find(',');
914-
if (comma != std::string::npos) first = first.substr(0, comma);
915-
auto type = common_speculative_type_from_name(first);
916-
if (type != COMMON_SPECULATIVE_TYPE_COUNT) {
917-
params.speculative.type = type;
918-
}
919-
#else
920900
// Upstream switched to a vector of types (comma-separated for multi-type
921901
// chaining via common_speculative_types_from_names). We keep accepting a
922902
// single value here, but also tolerate comma-separated lists.
@@ -945,7 +925,6 @@ static void params_parse(server_context& /*ctx_server*/, const backend::ModelOpt
945925
if (!parsed.empty()) {
946926
params.speculative.types = parsed;
947927
}
948-
#endif
949928
} else if (!strcmp(optname, "spec_n_max") || !strcmp(optname, "draft_max")) {
950929
if (optval != NULL) {
951930
try { params.speculative.draft.n_max = std::stoi(optval_str); } catch (...) {}
@@ -983,21 +962,6 @@ static void params_parse(server_context& /*ctx_server*/, const backend::ModelOpt
983962
// shares the target context size. Accept the option for backward
984963
// compatibility but silently ignore it.
985964

986-
// Everything below relies on struct shape introduced in ggml-org/llama.cpp#22838
987-
// (parallel drafting): `ngram_mod`, `ngram_map_k`, `ngram_map_k4v`,
988-
// `ngram_cache`, and the `draft.{cache_type_*, cpuparams*, tensor_buft_overrides}`
989-
// fields. The turboquant fork branched before that, so its build defines
990-
// LOCALAI_LEGACY_LLAMA_CPP_SPEC via patch-grpc-server.sh and these option
991-
// keys become unrecognized (silently dropped, like any unknown opt) for it.
992-
//
993-
// The `#ifdef LOCALAI_LEGACY_LLAMA_CPP_SPEC` / `#else` split below sits at the
994-
// closing-brace position of the `draft_ctx_size` branch on purpose: in the
995-
// legacy build the chain ends here (the brace closes draft_ctx_size), and in
996-
// the modern build the chain continues with `} else if (...)` instead, so the
997-
// brace count stays balanced under both branches of the preprocessor.
998-
#ifdef LOCALAI_LEGACY_LLAMA_CPP_SPEC
999-
}
1000-
#else
1001965
// --- ngram_mod family (upstream --spec-ngram-mod-*) ---
1002966
} else if (!strcmp(optname, "spec_ngram_mod_n_min")) {
1003967
if (optval != NULL) {
@@ -1127,7 +1091,6 @@ static void params_parse(server_context& /*ctx_server*/, const backend::ModelOpt
11271091
}
11281092
if (!cur.empty()) flush(cur);
11291093
}
1130-
#endif // LOCALAI_LEGACY_LLAMA_CPP_SPEC — closes the `else`/`#ifdef` opened at draft_ctx_size
11311094
}
11321095

11331096
// Set params.n_parallel from environment variable if not set via options (fallback)
@@ -1177,15 +1140,11 @@ static void params_parse(server_context& /*ctx_server*/, const backend::ModelOpt
11771140
params.tensor_buft_overrides.push_back({nullptr, nullptr});
11781141
}
11791142
}
1180-
// The draft tensor_buft_overrides are only populated under the modern
1181-
// (post-#22838) layout, whose population code is itself gated by
1182-
// LOCALAI_LEGACY_LLAMA_CPP_SPEC above. The turboquant fork lacks
1183-
// common_params_speculative::draft entirely, so skip the sentinel there too.
1184-
#ifndef LOCALAI_LEGACY_LLAMA_CPP_SPEC
1143+
// Terminate the draft tensor_buft_overrides list with a sentinel, mirroring
1144+
// the main-model handling above.
11851145
if (!params.speculative.draft.tensor_buft_overrides.empty()) {
11861146
params.speculative.draft.tensor_buft_overrides.push_back({nullptr, nullptr});
11871147
}
1188-
#endif
11891148

11901149
// TODO: Add yarn
11911150

backend/cpp/turboquant/Makefile

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11

22
# Pinned to the HEAD of feature/turboquant-kv-cache on https://github.com/TheTom/llama-cpp-turboquant.
33
# Auto-bumped nightly by .github/workflows/bump_deps.yaml.
4-
TURBOQUANT_VERSION?=5aeb2fdbe26cd4c534c6fa15de73cb5749bd0403
4+
TURBOQUANT_VERSION?=7d9715f1f071fa07c7b2ad3dbfd320b314139e65
55
LLAMA_REPO?=https://github.com/TheTom/llama-cpp-turboquant
66

77
CMAKE_ARGS?=

backend/cpp/turboquant/patch-grpc-server.sh

Lines changed: 26 additions & 80 deletions
Original file line numberDiff line numberDiff line change
@@ -4,21 +4,19 @@
44
#
55
# 1. Augment the kv_cache_types[] allow-list so `LoadModel` accepts the
66
# fork-specific `turbo2` / `turbo3` / `turbo4` cache types.
7-
# 2. Replace `get_media_marker()` (added upstream in ggml-org/llama.cpp#21962,
8-
# server-side random per-instance marker) with the legacy "<__media__>"
9-
# literal. The fork branched before that PR, so server-common.cpp has no
10-
# get_media_marker symbol. The fork's mtmd_default_marker() still returns
11-
# "<__media__>", and Go-side tooling falls back to that sentinel when the
12-
# backend does not expose media_marker, so substituting the literal keeps
13-
# behavior identical on the turboquant path.
14-
# 3. Revert the `common_params_speculative` field references to the
15-
# pre-refactor flat layout. Upstream ggml-org/llama.cpp#22397 split the
16-
# struct into nested `draft` / `ngram_simple` / `ngram_mod` / etc. members;
17-
# the turboquant fork branched before that PR and still exposes the flat
18-
# `n_max`, `mparams_dft`, `ngram_size_n`, ... fields. The substitutions
19-
# below map the new nested paths back to the legacy flat names so the
20-
# shared grpc-server.cpp keeps compiling against the fork's common.h.
21-
# Drop this block once the fork rebases past #22397.
7+
# 2. Define LOCALAI_TURBOQUANT_NO_CHECKPOINT_MIN_STEP at the top of the file
8+
# so the grpc-server option parser skips the two references to
9+
# common_params::checkpoint_min_step (the default and the option handler).
10+
# That field does not exist in the fork yet; drop this once it does.
11+
#
12+
# The fork used to lag upstream on the whole common_params_speculative refactor
13+
# (ggml-org/llama.cpp#22397/#22838/#22964), the model_tgt rename (#22838) and
14+
# get_media_marker (#21962), which required a much larger compat shim here
15+
# (flat-field sed renames + a coarse LOCALAI_LEGACY_LLAMA_CPP_SPEC define). The
16+
# fork has since rebased past all of those, so the only remaining gap is
17+
# checkpoint_min_step. If a future bump reintroduces a divergence, add a narrow
18+
# guard in grpc-server.cpp keyed on a fork-specific macro and inject it here
19+
# rather than resurrecting the coarse one.
2220
#
2321
# We patch the *copy* sitting in turboquant-<flavor>-build/, never the original
2422
# under backend/cpp/llama-cpp/, so the stock llama-cpp build keeps compiling
@@ -72,86 +70,34 @@ else
7270
echo "==> KV allow-list patch OK"
7371
fi
7472

75-
if grep -q 'get_media_marker()' "$SRC"; then
76-
echo "==> patching $SRC to replace get_media_marker() with legacy \"<__media__>\" literal"
77-
# Only one call site today (ModelMetadata), but replace all occurrences to
78-
# stay robust if upstream adds more. Use a temp file to avoid relying on
79-
# sed -i portability (the builder image uses GNU sed, but keeping this
80-
# consistent with the awk block above).
81-
sed 's/get_media_marker()/"<__media__>"/g' "$SRC" > "$SRC.tmp"
82-
mv "$SRC.tmp" "$SRC"
83-
echo "==> get_media_marker() substitution OK"
84-
else
85-
echo "==> $SRC has no get_media_marker() call, skipping media-marker patch"
86-
fi
87-
88-
if grep -q 'params\.speculative\.draft\.\|params\.speculative\.ngram_simple\.' "$SRC"; then
89-
echo "==> patching $SRC to revert common_params_speculative refs to pre-#22397 flat layout"
90-
# Each substitution is the exact post-refactor path → legacy flat field.
91-
# Order doesn't matter because the source paths are disjoint, but we keep
92-
# the most-specific (mparams.path) first for readability.
93-
sed -E \
94-
-e 's/params\.speculative\.draft\.mparams\.path/params.speculative.mparams_dft.path/g' \
95-
-e 's/params\.speculative\.draft\.n_max/params.speculative.n_max/g' \
96-
-e 's/params\.speculative\.draft\.n_min/params.speculative.n_min/g' \
97-
-e 's/params\.speculative\.draft\.p_min/params.speculative.p_min/g' \
98-
-e 's/params\.speculative\.draft\.p_split/params.speculative.p_split/g' \
99-
-e 's/params\.speculative\.draft\.n_gpu_layers/params.speculative.n_gpu_layers/g' \
100-
-e 's/params\.speculative\.draft\.n_ctx/params.speculative.n_ctx/g' \
101-
-e 's/params\.speculative\.ngram_simple\.size_n/params.speculative.ngram_size_n/g' \
102-
-e 's/params\.speculative\.ngram_simple\.size_m/params.speculative.ngram_size_m/g' \
103-
-e 's/params\.speculative\.ngram_simple\.min_hits/params.speculative.ngram_min_hits/g' \
104-
"$SRC" > "$SRC.tmp"
105-
mv "$SRC.tmp" "$SRC"
106-
echo "==> speculative field rename OK"
107-
else
108-
echo "==> $SRC has no post-#22397 speculative field refs, skipping spec rename patch"
109-
fi
110-
111-
# 4. Revert the `ctx_server.impl->model_tgt` rename introduced by upstream
112-
# ggml-org/llama.cpp#22838 (parallel drafting). The turboquant fork still
113-
# exposes the field as `model` on `server_context_impl`. The two call sites
114-
# are in the Rerank and ModelMetadata RPC handlers.
115-
if grep -q 'ctx_server\.impl->model_tgt' "$SRC"; then
116-
echo "==> patching $SRC to revert ctx_server.impl->model_tgt -> ctx_server.impl->model"
117-
sed -E 's/ctx_server\.impl->model_tgt/ctx_server.impl->model/g' "$SRC" > "$SRC.tmp"
118-
mv "$SRC.tmp" "$SRC"
119-
echo "==> model_tgt rename OK"
120-
else
121-
echo "==> $SRC has no ctx_server.impl->model_tgt refs, skipping model_tgt rename patch"
122-
fi
123-
124-
# 5. Define LOCALAI_LEGACY_LLAMA_CPP_SPEC at the top of the file so the
125-
# grpc-server option parser skips the new option-handler blocks (ngram_mod,
126-
# ngram_map_k, ngram_map_k4v, ngram_cache, draft.cache_type_*, draft.cpuparams*,
127-
# draft.tensor_buft_overrides) introduced for the post-#22838 layout, the
128-
# draft.tensor_buft_overrides sentinel termination, and the
129-
# common_params::checkpoint_min_step default/option (added with the
130-
# 35c9b1f3 bump). Those blocks reference struct fields that simply do not
131-
# exist in the fork.
132-
if grep -q '^#define LOCALAI_LEGACY_LLAMA_CPP_SPEC' "$SRC"; then
133-
echo "==> $SRC already defines LOCALAI_LEGACY_LLAMA_CPP_SPEC, skipping"
73+
# 2. Define LOCALAI_TURBOQUANT_NO_CHECKPOINT_MIN_STEP at the top of the file so
74+
# the grpc-server option parser skips the two references to
75+
# common_params::checkpoint_min_step (the default assignment and the option
76+
# handler). That field does not exist in the fork yet. Drop this block once
77+
# the fork rebases past the bump that added checkpoint_min_step.
78+
if grep -q '^#define LOCALAI_TURBOQUANT_NO_CHECKPOINT_MIN_STEP' "$SRC"; then
79+
echo "==> $SRC already defines LOCALAI_TURBOQUANT_NO_CHECKPOINT_MIN_STEP, skipping"
13480
else
135-
echo "==> patching $SRC to define LOCALAI_LEGACY_LLAMA_CPP_SPEC at the top"
136-
# Insert the define before the very first `#include` so it precedes all the
137-
# speculative-decoding code paths.
81+
echo "==> patching $SRC to define LOCALAI_TURBOQUANT_NO_CHECKPOINT_MIN_STEP at the top"
82+
# Insert the define before the very first `#include` so it precedes the
83+
# checkpoint_min_step references.
13884
awk '
13985
!done && /^#include/ {
140-
print "#define LOCALAI_LEGACY_LLAMA_CPP_SPEC 1"
86+
print "#define LOCALAI_TURBOQUANT_NO_CHECKPOINT_MIN_STEP 1"
14187
print "// ^ injected by backend/cpp/turboquant/patch-grpc-server.sh"
14288
print ""
14389
done = 1
14490
}
14591
{ print }
14692
END {
14793
if (!done) {
148-
print "patch-grpc-server.sh: no #include anchor found to insert LOCALAI_LEGACY_LLAMA_CPP_SPEC" > "/dev/stderr"
94+
print "patch-grpc-server.sh: no #include anchor found to insert LOCALAI_TURBOQUANT_NO_CHECKPOINT_MIN_STEP" > "/dev/stderr"
14995
exit 1
15096
}
15197
}
15298
' "$SRC" > "$SRC.tmp"
15399
mv "$SRC.tmp" "$SRC"
154-
echo "==> LOCALAI_LEGACY_LLAMA_CPP_SPEC define OK"
100+
echo "==> LOCALAI_TURBOQUANT_NO_CHECKPOINT_MIN_STEP define OK"
155101
fi
156102

157103
echo "==> all patches applied"
Lines changed: 55 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,55 @@
1+
hip: port the turboquant CUDA additions that ggml's HIP shim doesn't cover
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:
6+
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.
26+
27+
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
28+
index 0427e6b..6352e6a 100644
29+
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
30+
+++ 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_
48+
ggml_cuda_set_device(dev_ctx->device);
49+
50+
cudaEvent_t event;
51+
- CUDA_CHECK(cudaEventCreate(&event));
52+
+ CUDA_CHECK(cudaEventCreateWithFlags(&event, cudaEventDisableTiming));
53+
54+
return new ggml_backend_event {
55+
/* .device = */ dev,

0 commit comments

Comments
 (0)