[Captain-Tripps]

Overview

SYCL currently leaves set_tensor_2d unset, so uploads through ggml_backend_tensor_set_2d fall back to a separate set_tensor call for every row. In tensor-split mode (-sm tensor), that adds a lot of overhead when loading model weights.

This adds a SYCL implementation that packs the source rows into a host buffer and copies them in chunks of up to 64 MiB when the destination rows are contiguous. The limit keeps the temporary host allocation bounded. If the destination rows aren't contiguous, it uses the existing per-row copy path.

I tested this on two Arc Pro B70s with gpt-oss-120b UD-Q4_K_XL, comparing master f2918ca against the same commit with this change:

llama-perplexity -sm tensor -ngl 99 -c 512 --chunks 4 --load-mode dio

Across two runs, total runtime dropped from 263-268 seconds to 114-115 seconds. Earlier runs put the time to the first perplexity pass, which is mostly model loading, at roughly 256 seconds on master versus 110 seconds with the change.

For a direct copy check, I used a test build that read each uploaded chunk back and compared it with the source rows. One full model load checked 26,784,256 rows with no mismatches. Perplexity results were within the reported uncertainty.

The full test-backend-ops suite hits the same assertion in element_wise.cpp on both master and the patched build, just after the GEGLU_QUICK cases. That suite doesn't exercise set_tensor_2d; the readback comparison above checks the new upload path directly.

The branch was rebased onto f2cc728. There were no SYCL changes between that commit and this tested base, and the patch is unchanged.

Additional information

Requirements

• I have read and agree with the contributing guidelines - YES
• AI usage disclosure: YES - Claude helped me code/build/test this

[arthw]

It's good job!

The running time is reduced from 35s to 27s on 2xArc770 by the cmd:
time ./build/bin/llama-perplexity -sm tensor -ngl 99 -c 512 --chunks 4 --load-mode dio -m ../models/Meta-Llama-3.1-8B-Instruct-Q4_K_S.gguf -f README.md

Thank you!

[arthw]

In the first for(), call wait() for every memcpy() is not necessary.
Only need call wait() on once time out of the first for(), to reduce the performance impact.

[Captain-Tripps]

Thanks for reviewing and testing this on the A770s!

I updated the implementation to reuse the pinned staging slots already used by set_tensor. The original version reused a single buffer, so removing its' wait would let the CPU overwrite data while the device was still reading it.

With the staging slots, it only waits before reusing a slot, allowing packing and device copies to overlap.

On my two B70s with gpt-oss-120b, this reduced the time in set_tensor_2d from 17–18 seconds to 9–9.5 seconds, including waiting for the copies to finish. Total runtime dropped from 73–76 seconds to 64–67 seconds.

I also tested a single ext_oneapi_memcpy2d call, but it was considerably slower on this setup: around 227–229 seconds for the copies. The updated staging implementation passed a full readback comparison of 26.8 million rows with zero mismatches.

Please let me know if you have any other questions or requests. Hope this helps the community!

[Captain-Tripps]

Hi all - Anything else you need here? I have other PR's with additional improvements I'd like to submit when this one is approved or rejected.

[92es]

Tested with 2 - B70's loading a 35GB and 52GB LLM.

Deltas are repeatable across multiple test runs with minor variance.

llama-perplexity -m Qwen3.8-27B-UD-Q8_K_4XL.gguf -dev SYCL0,SYCL1 -sm tensor -ngl 99 -c 512 --chunks 4 --load-mode dio -v

35GB LLM	stock	patch	delta
load time	27639.29 ms	23893.03 ms	-13.55%
real	0m27.883s	0m24.131s	-13.46%
user	0m17.056s	0m13.440s	-21.20%
sys	0m3.850s	0m3.730s	-3.12%

llama-perplexity -m Qwen3.8-27B-BF16/Qwen3.8-27B-BF16-00001-of-00002.gguf -dev SYCL0,SYCL1 -sm tensor -ngl 99 -c 512 --chunks 4 --load-mode dio -v

52GB LLM	stock	patch	delta
load time	34443.25 ms	30314.14 ms	-11.99%
real	0m34.687s	0m30.549s	-11.93%
user	0m20.171s	0m15.938s	-20.99%
sys	0m4.717s	0m4.793s	+1.61%

Test System:
Debain Trixie Backports Kernel 7.2.9
Z790, i5-14600K, 96GB DDR5 6000 CL30, ramdisk
SYCL0 - Gen5x16 hung straight off CPU
SYCL1 - Gen4x4 hung under Z790's DMI Southbridge bus

@Captain-Tripps - Thanks much. Added to my daily driver.

[arthw]

It's good job!

I tested on Arc770, running time is decreased.

Thank you!

[moriz1]

i tested this with my 2x B70 and qwen3.8-flash-next IQ3_XS. -sm tensor now loads within a minute, as opposed to the 20+ minutes it took before. however performance was quite poor, much slower than -sm layer, and GPU power draw remained about the same. that said, performance with -sm tensor has always been poor on this model with my hardware; probably something else to look at next.

[92es]

@moriz1 -sm tensor favors token generation while -sm layer favors prompt eval with dual GPUs irrespective of this patch.

I measure a starting 300 token/s delta between file ingest between tensor=860 ingest/faster tg and layer=1160 ingest/slower tg.

This is a bus/memory speed issue between multiple GPUs ( and the code that does the data exchange ). I do not believe the B70 exposes direct P2P data exchanges.

[moriz1]

@92es by poor performance, I meant that token generation has gone DOWN with -sm tensor compared to -sm layer. 18 vs 26 t/s.

Like I alluded to earlier, this is likely a separate issue.

[92es]

@moriz1

I have just the other problem, but I am using a dense LLM across 2 B70s not the Flash-Next structure.

I am testing a small patch to ggml_sycl_comm_exchange() that brings split tensor ingest t/s on par with split layer ingest t/s. My tensor 860 t/s is now 1050 t/s, not quite the full 1160, but I closed 200:300 t/s gap between the 2.

This depends on your PCIe bus speeds. I have a Gen4x4 bottleneck that I am trying to make best use of the bandwidth.

Sneak Preview:

static void ggml_sycl_comm_exchange(ggml_backend_sycl_comm_context * comm, queue_ptr q0, queue_ptr q1,
                                    void * dst0, const void * src1, void * dst1, const void * src0,
                                    size_t nbytes) {

    // Reallocate host buffers if they are too small
    if (comm->host_bytes < nbytes) {
        q0->wait();
        q1->wait();
        comm->free_host();

        // OPTIMIZATION: Only 2 host buffers are required now instead of 4
        comm->host_out[0] = sycl::malloc_host<uint8_t>(nbytes, q0->get_context());
        comm->host_out[1] = sycl::malloc_host<uint8_t>(nbytes, q1->get_context());

        GGML_ASSERT(comm->host_out[0] && comm->host_out[1]);
        comm->host_bytes = nbytes;
    }

    // Step 1: Concurrently copy from GPUs into their respective host buffers (Asynchronous)
    sycl::event e_d2h_0 = q0->memcpy(comm->host_out[0], src0, nbytes);
    sycl::event e_d2h_1 = q1->memcpy(comm->host_out[1], src1, nbytes);

    // Step 2: Sync point. Wait until the host buffers are fully populated by the GPUs
    e_d2h_0.wait();
    e_d2h_1.wait();

    // Step 3: Cross-upload directly from the opposing host buffer.
    // CPU std::memcpy is eliminated. Both queues pull from host memory in parallel.
    sycl::event e_h2d_0 = q0->memcpy(dst0, comm->host_out[1], nbytes); // q0 pulls from host_out[1]
    sycl::event e_h2d_1 = q1->memcpy(dst1, comm->host_out[0], nbytes); // q1 pulls from host_out[0]

    // Step 4: Finalize the exchange before returning
    e_h2d_0.wait();
    e_h2d_1.wait();
}

Update: I unit tested the [4 buffer memcpy/wait/memcpy] pattern against the [2 buffer event/memccpy/wait] pattern before making the patch.

The Gen5x16 to Gen4x4 pathway went from 5.3 GB/s ( mainline code ) to 7.1 GB/s ( +34%, patch code ), almost maxing out the theoretical limit of 7.8.GB/s.

[92es]

@Captain-Tripps , can you check out #30956 ?

I do not want it to step on your #29604 quad B70 config.

[Captain-Tripps]

@Captain-Tripps , can you check out #30956 ?

I do not want it to step on your #29604 quad B70 config.

@92es - I tested #30956 on Intel Arc Pro B70 (SYCL, Windows 11, 2 cards via ONEAPI_DEVICE_SELECTOR=level_zero:0,1). Caveat: this was a hand port onto our own SYCL branch (cherry-pick conflicts), not this branch as-is. Only ggml-sycl.dll differed between the baseline and the PR build. Overall - Your PR is a clear win for tensor-split prompt processing on SYCL.

Details:
Correctness: llama-perplexity at ub 4 gives bit-identical results with and without the PR in -sm tensor:

Qwen3.8-27B Q8_0: 4.8336
Qwen3.8-27B Q6_K_XL: 4.8264

At larger batch sizes the values stay within the baseline’s own run-to-run spread.

Speed, 2-card -sm tensor, baseline → PR (server, about 33k-token prompt, A/B/B/A order, mean of 2 runs):

model prompt read decode
Mistral-Small-3.2-24B Q8_0 898 → 1024 tok/s (+14%) 21.3 → 21.8 (+2%)
gpt-oss-120b Q4_K_XL 454 → 478 tok/s (+5%) equal

Qwen3.8-27B, layer vs tensor+PR (131k context, q4_0 KV, MTP drafting with --spec-draft-n-max 4):

model split decode (MTP) 33k prompt read decode after it
Q8_0 layer 42.9 948 28.1
Q8_0 tensor + PR 35.1 (−18%) 941 (−1%) 22.3 (−21%)
Q6_K_XL layer 42.8 893 27.7
Q6_K_XL tensor + PR 35.1 (−18%) 878 (−2%) 24.7 (−11%)

This matches what you measured: Qwen3.8-27B prompt reads in tensor split land roughly level with layer split. With MTP drafting, though, tensor split decodes about 18% slower on 2 cards. I didn’t run tensor without the PR for this model, so I can’t say how much of that gap the PR closes.

Layer split + MTP is still the winning recipe on my machine and we've yet to unpack why. My guess is the per-layer exchange on each small draft/verify step, but that’s unmeasured.

[92es]

@Captain-Tripps , Is there a stable way to test MTP in llama-bench or llama-perplexity ( I am very knew at llama.cpp, but have coded and tested once or twice before ) ?

When I A/B test in llama-server with Twin B70s, MTP, 292K KV BF16 context having it read the same largest source code file on startup, ingest jumps from 860 - 1100 so ingest speedup is still present with MTP.

What is the a stable repeatable way to test MTP after the ingest ? When it is doing subsequent thinking, the modified pathway is still being called.

If I have a stable way to A/B the MTP, I can chase after the cause of the delta.

Thanks much.

[Captain-Tripps]

@92es
I don’t think llama-bench or llama-perplexity exercise the speculative/MTP path today, so they can’t give an end-to-end MTP number by themselves.

I’d split this into two measurements:

1. Measure the target verify-step cost with llama-bench, which should be very repeatable:
   llama-bench -m <model> -sm layer,tensor -fa 1 -p 1,2,5,8 -n 0 -d 0,32768 -r 10
   With --spec-draft-n-max 4, an MTP verify is one target forward over 5 tokens, so pp5 is the directly comparable shape. -d 0,32768 also covers the “after ingest” case. For pp5, verify-step time is 5000 / tok/s in ms.

2. Then measure actual MTP with a fixed prompt and greedy decoding:

llama-speculative-simple -m <model> -sm tensor \
  --spec-type draft-mtp --spec-draft-n-max 4 \
  -f prompt.txt -n 512 --temp 0 -c <ctx>

Repeat with -sm layer. This prints decoded tok/s plus n_drafted, n_accept, and acceptance %. The fixed prompt + --temp 0 should make it deterministic for a given build and split mode.

The important bit is to compare acceptance alongside tok/s. Tensor splitting changes reduction order slightly, which can change logits and therefore MTP acceptance. So an end-to-end tok/s difference can be either verify-step cost, acceptance, or both.

This should also be a useful check for the two-buffer change: compare p=1 and p=5 first. My current suspicion—based on code reading, not measurement—is that decode is dominated by the host waits in each allreduce, rather than the copy bandwidth. If that is right, the patch may help prompt throughput more than decode/MTP. I have not run the speculative-simple side on Zeus yet, so I’d treat that last part as a hypothesis to validate.

[92es]

@Captain-Tripps,

I think I may have a fix. Need to test more, but acceptance is now identical (mainline vs patch) down to 4 levels using single-shot llama-cli with MTP runs.

(0.862, 0.690, 0.552, 0.517) mean len = 3.62

Maybe able to have cake and eat it too !!!

I changed the (FIRST PAIR ONLY, not last pair) waits to ext_oneapi_submit_barrier to change how blocking behavior is done:

Before:

    // Step 2:
    e_d2h_0.wait();
    e_d2h_1.wait();

After:

    // Step 2:
    q0->ext_oneapi_submit_barrier({e_d2h_1}); // q0 waits for q1's device-to-host copy
    q1->ext_oneapi_submit_barrier({e_d2h_0}); // q1 waits for q0's device-to-host copy

Using:

   --spec-type draft-mtp \
   --spec-draft-n-max 4 \
   --spec-draft-p-min 0.75 \

Will be rerunning my entire llama-perplexity regression matrix in addition and updating #30956 accordingly.

Stock Run:

draft acceptance = 0.85393 (   76 accepted /    89 generated), mean len =  3.62
acc per pos = (0.862, 0.690, 0.552, 0.517)
draft-mtp: #calls(b,g,a) =    1     51     29, #gen drafts =     29, #acc drafts =    25, #gen tokens =     89, #acc tokens =    76, #mean acc len = 3.62, #acc rate/pos = (0.862, 0.690, 0.552, 0.517), dur(b,g,a) = 0.005, 362.976, 0.050 ms

Patch Run:

draft acceptance = 0.85393 (   76 accepted /    89 generated), mean len =  3.62
acc per pos = (0.862, 0.690, 0.552, 0.517)
draft-mtp: #calls(b,g,a) =    1     51     29, #gen drafts =     29, #acc drafts =    25, #gen tokens =     89, #acc tokens =    76, #mean acc len = 3.62, #acc rate/pos = (0.862, 0.690, 0.552, 0.517), dur(b,g,a) = 0.006, 366.711, 0.040 ms

Update:

• Fixed code above.
• Tested --spec-draft-n-max 6, all 6 positions are identical (mainline:patch)

[Captain-Tripps]

@92es - Nice. Will be interested to see how it goes. Let's move this collab out of this PR so @arthw can get this merged with main and we don't clutter it up. I have a public repo where I'm updating my hill-climbing efforts on B70's. That's where my PR's here stem from. https://github.com/Captain-Tripps/b70-lab/blob/main/llama.cpp-sycl/TENSOR-SPLIT-DECODE.md is the one I'm looking at now for tensor split. Point your clanker at https://github.com/Captain-Tripps/b70-lab and lets see what we can come up with.