Commit dbeb37548 for llama.cpp
commit dbeb37548e25abc6e54961c4c99e63f191367809
Author: Titaniumtown <titaniumtown@proton.me>
Date: Mon Sep 7 06:24:14 2026 -0700
sycl: add a batched L2_NORM kernel (#28222)
* sycl: add a batched L2_NORM kernel
* sycl: batch consecutive L2_NORM siblings in the graph dispatch
Measured on Intel Arc Pro B70 (Battlemage), Qwen3.6-27B Q4_K_M, f16 KV,
npp=128 ntg=128 npl=2, GGML_SYCL profiler:
L2_NORM dispatches 12480 -> 6240
L2_NORM device time 68.77 -> 39.14 ms (-43%)
total device time 6782 -> 6748 ms (-0.5%)
wall decode t/s flat
* tests: add L2_NORM_BATCH coverage
diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index bfe6f1016..4091f73a4 100644
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -4858,6 +4858,78 @@ static bool ggml_sycl_mul_mat_glu_mmvq_fused(ggml_backend_sycl_context & ctx, gg
/*stride_col_dst=*/(int) glu->ne[0], stream);
}
+// Batch the run of consecutive L2_NORM siblings starting at node_idx into one launch.
+// Returns the number of extra graph nodes consumed, or 0 if the run is shorter than two
+// (the caller then runs the norm through the per-tensor kernel).
+static int ggml_sycl_l2_norm_batch_fused(ggml_backend_sycl_context & ctx, ggml_cgraph * cgraph, int node_idx) {
+ const ggml_tensor * node = cgraph->nodes[node_idx];
+ if (ggml_sycl_info().device_count != 1 || node->type != GGML_TYPE_F32 ||
+ node->src[0]->type != GGML_TYPE_F32 || node->src[0]->ne[0] >= 1024) {
+ return 0;
+ }
+
+ ggml_tensor * batch[GGML_SYCL_L2_BATCH_MAX];
+ int count = 0;
+ int last = node_idx;
+ float eps0;
+ memcpy(&eps0, node->op_params, sizeof(float));
+
+ // Conservative aliasing test: the batched norms run concurrently in one kernel,
+ // so none may read what another writes, and none may write where another writes.
+ auto overlaps = [](const ggml_tensor * a, const ggml_tensor * b) {
+ const char * ab = (const char *) a->data;
+ const char * bb = (const char *) b->data;
+ return ab < bb + ggml_nbytes(b) && bb < ab + ggml_nbytes(a);
+ };
+
+ for (int j = node_idx; j < cgraph->n_nodes && count < GGML_SYCL_L2_BATCH_MAX; ++j) {
+ ggml_tensor * nj = cgraph->nodes[j];
+ if (ggml_is_empty(nj) || nj->op == GGML_OP_RESHAPE || nj->op == GGML_OP_TRANSPOSE ||
+ nj->op == GGML_OP_VIEW || nj->op == GGML_OP_PERMUTE || nj->op == GGML_OP_NONE ||
+ (nj->flags & GGML_TENSOR_FLAG_COMPUTE) == 0) {
+ continue; // not a launch; cannot break a run of adjacent norms
+ }
+ if (nj->op != GGML_OP_L2_NORM || nj->type != GGML_TYPE_F32 ||
+ nj->src[0]->type != GGML_TYPE_F32 || !ggml_are_same_shape(nj, node) ||
+ !ggml_are_same_shape(nj->src[0], node->src[0])) {
+ break; // any other launch ends the run
+ }
+ bool same_nb = true;
+ for (int d = 0; d < GGML_MAX_DIMS; ++d) {
+ if (nj->nb[d] != node->nb[d] || nj->src[0]->nb[d] != node->src[0]->nb[d]) {
+ same_nb = false;
+ break;
+ }
+ }
+ if (!same_nb) {
+ break; // one nb[] stride set is shared by the whole batch
+ }
+ float epsj;
+ memcpy(&epsj, nj->op_params, sizeof(float));
+ if (epsj != eps0) {
+ break; // eps mismatch ends the run
+ }
+ bool indep = true;
+ for (int k = 0; k < count; ++k) {
+ if (overlaps(nj->src[0], batch[k]) || overlaps(nj, batch[k])) {
+ indep = false;
+ break;
+ }
+ }
+ if (!indep) {
+ break; // an overlapping tensor would race inside one launch
+ }
+ batch[count++] = nj;
+ last = j;
+ }
+ if (count < 2) {
+ return 0; // a lone norm falls through to the per-tensor kernel
+ }
+ ggml_sycl_l2_norm_batch(ctx, batch, count);
+ return last - node_idx;
+}
+
+
__dpct_inline__ static void k_copy_src1_to_contiguous(
const char *__restrict__ src1_original, char *__restrict__ src1_contiguous,
const mmid_row_mapping *__restrict__ row_mapping,
@@ -5908,6 +5980,17 @@ static void ggml_backend_sycl_graph_compute_impl(ggml_backend_sycl_context * syc
continue;
}
+ // Batch consecutive independent same-shape F32 L2_NORM siblings (the GDN q/k
+ // norms) into one launch; sources are strided views of the fused qkv buffer, so
+ // the scan skips the interleaved view nodes instead of breaking on them.
+ if (node->op == GGML_OP_L2_NORM) {
+ const int l2_batch_skip = ggml_sycl_l2_norm_batch_fused(*sycl_ctx, cgraph, i);
+ if (l2_batch_skip > 0) {
+ i += l2_batch_skip;
+ continue;
+ }
+ }
+
if (node->op == GGML_OP_MUL_MAT && ggml_sycl_mul_mat_glu_mmvq_fused(*sycl_ctx, cgraph, i)) {
i += 2;
continue;
diff --git a/ggml/src/ggml-sycl/norm.cpp b/ggml/src/ggml-sycl/norm.cpp
index 2d3033729..bc36a9d4c 100644
--- a/ggml/src/ggml-sycl/norm.cpp
+++ b/ggml/src/ggml-sycl/norm.cpp
@@ -543,6 +543,62 @@ static void l2_norm_f32_sycl(const float * x,
}
}
+// Batched L2 norm: N independent same-shape F32 tensors in one launch; the tensor
+// index is folded into grid dim0 and each row's reduction is identical to the
+// single-tensor kernel, so the result is bit-exact.
+struct l2_batch_ptrs {
+ const float * src[GGML_SYCL_L2_BATCH_MAX];
+ float * dst[GGML_SYCL_L2_BATCH_MAX];
+};
+
+// One stride set shared by the whole batch: the caller only groups tensors whose nb[]
+// all match, so per-tensor state stays two pointers.
+struct l2_batch_strides {
+ int ne1, ne2;
+ int64_t ss0, ss1, ss2, ss3;
+ int64_t ds0, ds1, ds2, ds3;
+};
+
+template <int warp_size>
+static void l2_norm_f32_batch(l2_batch_ptrs p, l2_batch_strides st, const int ncols, const float eps,
+ const sycl::nd_item<3> & item_ct1) {
+ const int t = item_ct1.get_group(0); // tensor index
+ const int r = item_ct1.get_group(2); // flattened row over ne1*ne2*ne3
+ const int tid = item_ct1.get_local_id(2);
+
+ const int i1 = r % st.ne1;
+ const int i2 = (r / st.ne1) % st.ne2;
+ const int i3 = r / (st.ne1 * st.ne2);
+
+ const float * x = p.src[t] + i3 * st.ss3 + i2 * st.ss2 + i1 * st.ss1;
+ float * dst = p.dst[t] + i3 * st.ds3 + i2 * st.ds2 + i1 * st.ds1;
+
+ float tmp = 0.0f;
+ for (int col = tid; col < ncols; col += warp_size) {
+ const float xi = x[col * st.ss0];
+ tmp += xi * xi;
+ }
+ tmp = block_reduce<block_reduce_method::SUM, warp_size>(tmp, (float *) nullptr, warp_size);
+ const float scale = sycl::rsqrt(sycl::fmax(tmp, eps * eps));
+ for (int col = tid; col < ncols; col += warp_size) {
+ dst[col * st.ds0] = scale * x[col * st.ss0];
+ }
+}
+
+template <int warp_size>
+static void l2_norm_f32_batch_sycl(l2_batch_ptrs p, l2_batch_strides st, const int n_tensors,
+ const int ncols, const int nrows_total, const float eps,
+ queue_ptr stream) {
+ const dpct::dim3 blocks_num(nrows_total, 1, n_tensors);
+ const dpct::dim3 block_dims(warp_size, 1, 1);
+ stream->submit([&](sycl::handler & cgh) {
+ cgh.parallel_for(sycl::nd_range<3>(blocks_num * block_dims, block_dims),
+ [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(warp_size)]] {
+ l2_norm_f32_batch<warp_size>(p, st, ncols, eps, item_ct1);
+ });
+ });
+}
+
void ggml_sycl_op_norm(ggml_backend_sycl_context& ctx, ggml_tensor* dst) {
const ggml_tensor * src0 = dst->src[0];
@@ -961,3 +1017,30 @@ void ggml_sycl_op_l2_norm(ggml_backend_sycl_context& ctx, ggml_tensor* dst) {
l2_norm_f32_sycl<WARP_SIZE>(src0_d, dst_d, ne00, ne01, ne02, ne03,
ss0, ss1, ss2, ss3, ds0, ds1, ds2, ds3, eps, stream, ctx.device);
}
+
+// nodes[0..count) are independent, same-shape, same-eps, same-nb L2_NORM ops validated
+// by the caller; requires ncols < 1024 (the warp reduction path).
+void ggml_sycl_l2_norm_batch(ggml_backend_sycl_context & ctx, ggml_tensor ** nodes, int count) {
+ const ggml_tensor * s0 = nodes[0]->src[0];
+ const int ncols = (int) s0->ne[0];
+ const int nrows_total = (int) ggml_nrows(s0);
+ float eps;
+ memcpy(&eps, nodes[0]->op_params, sizeof(float));
+ GGML_ASSERT(eps >= 0.0f);
+
+ l2_batch_ptrs p{};
+ for (int t = 0; t < count; ++t) {
+ p.src[t] = (const float *) nodes[t]->src[0]->data;
+ p.dst[t] = (float *) nodes[t]->data;
+ }
+
+ const ggml_tensor * d0 = nodes[0];
+ const size_t ts = ggml_type_size(GGML_TYPE_F32);
+ l2_batch_strides st{};
+ st.ne1 = (int) s0->ne[1];
+ st.ne2 = (int) s0->ne[2];
+ st.ss0 = s0->nb[0] / ts; st.ss1 = s0->nb[1] / ts; st.ss2 = s0->nb[2] / ts; st.ss3 = s0->nb[3] / ts;
+ st.ds0 = d0->nb[0] / ts; st.ds1 = d0->nb[1] / ts; st.ds2 = d0->nb[2] / ts; st.ds3 = d0->nb[3] / ts;
+
+ l2_norm_f32_batch_sycl<WARP_SIZE>(p, st, count, ncols, nrows_total, eps, ctx.stream());
+}
diff --git a/ggml/src/ggml-sycl/norm.hpp b/ggml/src/ggml-sycl/norm.hpp
index ef7b2d386..46c6de2a1 100644
--- a/ggml/src/ggml-sycl/norm.hpp
+++ b/ggml/src/ggml-sycl/norm.hpp
@@ -29,4 +29,7 @@ void ggml_sycl_op_group_norm(ggml_backend_sycl_context& ctx, ggml_tensor* dst);
void ggml_sycl_op_l2_norm(ggml_backend_sycl_context& ctx, ggml_tensor* dst);
+#define GGML_SYCL_L2_BATCH_MAX 8
+void ggml_sycl_l2_norm_batch(ggml_backend_sycl_context & ctx, ggml_tensor ** nodes, int count);
+
#endif // GGML_SYCL_NORM_HPP
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index 19eaacbec..cd16e5851 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -7206,6 +7206,49 @@ struct test_group_norm_mul_add : public test_case {
}
};
+// GGML_OP_L2_NORM x N: independent same-shape norms in one graph (strided qkv views or
+// contiguous), consuming adds nested so the norms stay adjacent in the graph.
+struct test_l2_norm_batch : public test_case {
+ const ggml_type type;
+ const std::array<int64_t, 4> ne;
+ const int n_norms;
+ const float eps;
+ const bool strided;
+
+ std::string vars() override { return VARS_TO_STR5(type, ne, n_norms, eps, strided); }
+ std::string op_desc(ggml_tensor * t) override { GGML_UNUSED(t); return "L2_NORM_BATCH"; }
+ bool run_whole_graph() override { return true; }
+
+ test_l2_norm_batch(ggml_type type = GGML_TYPE_F32, std::array<int64_t, 4> ne = { 128, 16, 16, 1 },
+ int n_norms = 4, float eps = 1e-12f, bool strided = true)
+ : type(type), ne(ne), n_norms(n_norms), eps(eps), strided(strided) {}
+
+ ggml_tensor * build_graph(ggml_context * ctx) override {
+ GGML_ASSERT(n_norms >= 2 && n_norms <= 8);
+ ggml_tensor * parent = nullptr;
+ if (strided) {
+ parent = ggml_new_tensor_4d(ctx, type, ne[0], ne[1] * n_norms, ne[2], ne[3]); // qkv buffer
+ }
+ ggml_tensor * norms[8];
+ for (int t = 0; t < n_norms; ++t) {
+ ggml_tensor * src;
+ if (strided) {
+ src = ggml_view_4d(ctx, parent, ne[0], ne[1], ne[2], ne[3], parent->nb[1], parent->nb[2],
+ parent->nb[3], t * ne[1] * parent->nb[1]);
+ } else {
+ src = ggml_new_tensor(ctx, type, 4, ne.data());
+ }
+ norms[t] = ggml_l2_norm(ctx, src, eps);
+ }
+ ggml_tensor * out = norms[n_norms - 1];
+ for (int t = n_norms - 2; t >= 0; --t) {
+ out = ggml_add(ctx, norms[t], out);
+ }
+ ggml_set_name(out, "out");
+ return out;
+ }
+};
+
// GGML_OP_L2_NORM
struct test_l2_norm : public test_case {
const ggml_type type;
@@ -9495,6 +9538,10 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
test_cases.emplace_back(new test_l2_norm(GGML_TYPE_F32, { n, 5, 4, 3 }, eps, false));
test_cases.emplace_back(new test_l2_norm(GGML_TYPE_F32, { n, 5, 4, 3 }, eps, true));
test_cases.emplace_back(new test_l2_norm(GGML_TYPE_F32, { n, 5, 4, 3 }, eps, false, true));
+ // sibling batching: strided (production shape) and contiguous, 2 and 4 wide
+ test_cases.emplace_back(new test_l2_norm_batch(GGML_TYPE_F32, { n, 5, 4, 3 }, 2, eps, true));
+ test_cases.emplace_back(new test_l2_norm_batch(GGML_TYPE_F32, { n, 5, 4, 3 }, 4, eps, true));
+ test_cases.emplace_back(new test_l2_norm_batch(GGML_TYPE_F32, { n, 5, 4, 3 }, 4, eps, false));
}
// row lengths that are not a multiple of 32, for the scalar (33) and float4 (132, 260) paths
for (uint32_t n : { 33, 132, 260 }) {
@@ -11181,6 +11228,16 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
}
}
+ // launch-overhead isolation: single L2_NORM launch vs batched siblings at the GDN
+ // production shape (strided qkv views) -- perf-mode only, the eval list has its own
+ // 2/4-wide coverage
+ for (int n : { 128, 256 }) {
+ test_cases.emplace_back(new test_l2_norm(GGML_TYPE_F32, { n, 16, 16, 1 }, 1e-12f, false, false));
+ test_cases.emplace_back(new test_l2_norm_batch(GGML_TYPE_F32, { n, 16, 16, 1 }, 2, 1e-12f, true));
+ test_cases.emplace_back(new test_l2_norm_batch(GGML_TYPE_F32, { n, 16, 16, 1 }, 4, 1e-12f, true));
+ }
+
+
return test_cases;
}