From b61186de014210b7c8dd10ab420d48635d73400e Mon Sep 17 00:00:00 2001 From: Titaniumtown Date: Mon, 7 Sep 2026 06:24:14 -0700 Subject: [PATCH] sycl: add a batched L2_NORM kernel (llama/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 --- ggml/src/ggml-sycl/ggml-sycl.cpp | 83 ++++++++++++++++++++++++++++++++ ggml/src/ggml-sycl/norm.cpp | 83 ++++++++++++++++++++++++++++++++ ggml/src/ggml-sycl/norm.hpp | 3 ++ 3 files changed, 169 insertions(+) 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 +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(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 +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(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(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(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