Commit 02269b4c0 for llama.cpp
commit 02269b4c00ec5b659f481e85178f9a167460fad5
Author: Captain-Tripps <jstaples2@hotmail.com>
Date: Mon Oct 26 09:12:25 2026 -0500
sycl: add set_tensor_2d for faster tensor-split loading (#30358)
* sycl: add set_tensor_2d for faster tensor-split loading
Pack host rows into chunks of up to 64 MiB when destination rows are
contiguous, avoiding a separate upload call for each row. Keep the
existing per-row copy path for non-contiguous destinations.
On two Arc Pro B70s with gpt-oss-120b UD-Q4_K_XL, total llama-perplexity
runtime drops from 263-268 s to 114-115 s. A full-load readback check
verified 26,784,256 rows with no mismatches.
* sycl: batch 2D tensor uploads through shared staging slots
Implement set_tensor_2d to speed up tensor-split model loading by
packing rows into the existing pinned staging slots and uploading
each batch with a single device copy.
Share staging initialization with set_tensor and wait only when
reusing a slot, allowing CPU packing and device copies to overlap.
Fall back to per-row uploads for noncontiguous device rows,
oversized rows, or unavailable staging.
---------
Co-authored-by: Jon Staples <250234834+Captain-Tripps@users.noreply.github.com>
diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index a04470116..ae0f291c4 100644
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -842,6 +842,19 @@ catch (sycl::exception const &exc) {
std::exit(1);
}
+// allocates the pinned upload staging on first use; returns false if staging is disabled or the allocation failed
+static bool ggml_sycl_buffer_init_staging(ggml_backend_sycl_buffer_context * ctx) {
+ const int n_slots = g_ggml_sycl_upload_staging_slots;
+ if (n_slots > 0 && ctx->staging.data == nullptr) {
+ ctx->staging.data = sycl::malloc_host(n_slots * ctx->staging_slot_size, *ctx->stream);
+ if (ctx->staging.data != nullptr) {
+ ctx->staging.events.resize(n_slots);
+ ctx->staging.submitted.assign(n_slots, false);
+ }
+ }
+ return ctx->staging.data != nullptr;
+}
+
static void ggml_backend_sycl_buffer_set_tensor(ggml_backend_buffer_t buffer,
ggml_tensor *tensor,
const void *data, size_t offset,
@@ -854,15 +867,7 @@ static void ggml_backend_sycl_buffer_set_tensor(ggml_backend_buffer_t buffer,
// copy through pinned memory so the device never reads mmap()ed pages directly
// chunks pipeline on the in-order compute queue, so no drain per tensor is needed
- const int n_slots = g_ggml_sycl_upload_staging_slots;
- if (n_slots > 0 && ctx->staging.data == nullptr) {
- ctx->staging.data = sycl::malloc_host(n_slots * ctx->staging_slot_size, *ctx->stream);
- if (ctx->staging.data != nullptr) {
- ctx->staging.events.resize(n_slots);
- ctx->staging.submitted.assign(n_slots, false);
- }
- }
- if (ctx->staging.data != nullptr) {
+ if (ggml_sycl_buffer_init_staging(ctx)) {
queue_ptr stream = ctx->stream;
char * dst = (char *) tensor->data + offset;
const char * src = (const char *) data;
@@ -904,6 +909,49 @@ catch (sycl::exception const &exc) {
std::exit(1);
}
+// Copy n_copies rows in one call. -sm tensor uploads split weights this way; the generic fallback does one set_tensor per row.
+static void ggml_backend_sycl_buffer_set_tensor_2d(ggml_backend_buffer_t buffer, ggml_tensor * tensor,
+ const void * data, size_t offset, size_t size,
+ size_t n_copies, size_t stride_tensor,
+ size_t stride_data) try {
+ ggml_backend_sycl_buffer_context * ctx = (ggml_backend_sycl_buffer_context *) buffer->context;
+ const char * src = (const char *) data;
+ ggml_sycl_set_device(ctx->device);
+
+ if (stride_tensor != size || size > ctx->staging_slot_size || !ggml_sycl_buffer_init_staging(ctx)) {
+ // dst rows not contiguous, row larger than a staging slot, or no staging: same as the generic fallback
+ for (size_t i = 0; i < n_copies; ++i) {
+ ggml_backend_sycl_buffer_set_tensor(buffer, tensor, src + i * stride_data, offset + i * stride_tensor, size);
+ }
+ return;
+ }
+
+ // dst rows are contiguous: pack as many rows as fit into a staging slot and copy each slot with one memcpy
+ // a slot is only waited on when it is reused, so packing the next slot overlaps with the previous copies
+ queue_ptr stream = ctx->stream;
+ char * dst = (char *) tensor->data + offset;
+ const size_t rows_per_slot = ctx->staging_slot_size / size;
+ for (size_t i0 = 0; i0 < n_copies; i0 += rows_per_slot) {
+ const size_t n = std::min(rows_per_slot, n_copies - i0);
+ const int slot = ctx->staging.next;
+ ctx->staging.next = (ctx->staging.next + 1) % (int) ctx->staging.submitted.size();
+ if (ctx->staging.submitted[slot]) {
+ ctx->staging.events[slot].wait_and_throw();
+ }
+ char * stage = (char *) ctx->staging.data + slot * ctx->staging_slot_size;
+ for (size_t i = 0; i < n; ++i) {
+ memcpy(stage + i * size, src + (i0 + i) * stride_data, size);
+ }
+ ctx->staging.events[slot] = stream->memcpy(dst + i0 * size, stage, n * size);
+ ctx->staging.submitted[slot] = true;
+ }
+}
+catch (sycl::exception const &exc) {
+ std::cerr << exc.what() << "Exception caught at file:" << __FILE__
+ << ", line:" << __LINE__ << std::endl;
+ std::exit(1);
+}
+
static void ggml_backend_sycl_buffer_get_tensor(ggml_backend_buffer_t buffer,
const ggml_tensor *tensor,
void *data, size_t offset,
@@ -1122,7 +1170,7 @@ static const ggml_backend_buffer_i ggml_backend_sycl_buffer_interface = {
/* .memset_tensor = */ ggml_backend_sycl_buffer_memset_tensor,
/* .set_tensor = */ ggml_backend_sycl_buffer_set_tensor,
/* .get_tensor = */ ggml_backend_sycl_buffer_get_tensor,
- /* .set_tensor_2d = */ NULL,
+ /* .set_tensor_2d = */ ggml_backend_sycl_buffer_set_tensor_2d,
/* .get_tensor_2d = */ NULL,
/* .cpy_tensor = */ ggml_backend_sycl_buffer_cpy_tensor,
/* .clear = */ ggml_backend_sycl_buffer_clear,