Lumpiasty/llama.cpp
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

sycl: fuse UNARY(silu|sigmoid|softplus) + MUL (#26411)

Browse Source 
Measured on Arc Pro B70 (Battlemage), Qwen3.6-27B Q4_K_M, -fa on, f16 KV,
-b 2048 -ub 2048, llama-bench -r 3, three interleaved A/B rounds:

  pp2048        1014.70 -> 1018.56 t/s   (+0.38%, within run-to-run spread)
  tg128         23.73 -> 23.86 t/s       (+0.57%)
  tg128 @ d4096 22.71 -> 22.86 t/s       (+0.62%)
This commit is contained in:
Titaniumtown
2026-08-13 11:41:38 +03:00
committed by GitHub
parent 8efbf65dbd
commit 1ee1cd9bc6
6 changed files with 280 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-sycl/element_wise.cpp
+87
View File
@@ -448,6 +448,47 @@ static void unary_gated_op_generic_kernel(
}
}
// Fused UNARY + MUL. Unlike the gated ops above, `x` and `g` are separate tensors of the
// same shape; `o0`/`o1` are their row strides in elements, so a half-view needs no repack.
// `dst` is contiguous and indexed flat. Math is done in f32, as the CPU and CUDA references do.
template<typename T, typename F>
static void unary_mul_flat_kernel(const T * x, const T * g, T * dst, const int64_t k, const sycl::nd_item<1> &item_ct1, F op) {
SYCL_GLOBAL_ID_LOOP(k, item_ct1) {
dst[i] = (T) (op((float) x[i]) * (float) g[i]);
}
}
template<typename T, typename F>
static void unary_mul_strided_kernel(const T * x, const T * g, T * dst, const int64_t k, const sycl::uint3 n_fd, const int64_t o0, const int64_t o1, const sycl::nd_item<1> &item_ct1, F op) {
SYCL_GLOBAL_ID_LOOP(k, item_ct1) {
const sycl::uint2 rc = fast_div_modulo((uint32_t) i, n_fd);
const int64_t j0 = rc.x() * o0 + rc.y();
const int64_t j1 = o0 == o1 ? j0 : rc.x() * o1 + rc.y();
dst[i] = (T) (op((float) x[j0]) * (float) g[j1]);
}
}
template<typename T, typename F>
static void unary_mul_sycl(const T * x, const T * g, T * dst, const int64_t k, const int64_t n, const int64_t o0, const int64_t o1, queue_ptr main_stream, F op) {
const size_t num_blocks = ceil_div((size_t) k, (size_t) SYCL_GLU_BLOCK_SIZE);
const sycl::nd_range<1> range(num_blocks * sycl::range<1>(SYCL_GLU_BLOCK_SIZE), sycl::range<1>(SYCL_GLU_BLOCK_SIZE));
// o0 == o1 == n makes (i/n)*o0 + (i%n) == i, so the strided kernel degenerates to the flat one
if (o0 == n && o1 == n) {
main_stream->parallel_for(range, [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
unary_mul_flat_kernel(x, g, dst, k, item_ct1, op);
});
return;
}
// 32-bit fastdiv, exact only below 2^31; ggml_sycl_can_fuse() already declined past that
GGML_ASSERT(k < ((int64_t) 1 << 31));
const sycl::uint3 n_fd = init_fastdiv_values((uint32_t) n);
main_stream->parallel_for(range, [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
unary_mul_strided_kernel(x, g, dst, k, n_fd, o0, o1, item_ct1, op);
});
}
namespace ggml_sycl_detail {
static void acc_f32_sycl(const char *x, const char *y, float *dst,
const int64_t n_elements,
@@ -991,6 +1032,52 @@ static inline void ggml_sycl_op_swiglu(ggml_backend_sycl_context & ctx, ggml_ten
});
}
// dst = op(unary_node->src[0]) * other, written straight to the MUL output, saving the
// standalone unary launch. Preconditions come from ggml_sycl_can_fuse(); re-asserted here.
void ggml_sycl_op_unary_mul_fused(ggml_backend_sycl_context & ctx, ggml_tensor * unary_node, ggml_tensor * mul_node) {
scope_op_debug_print scope_dbg_print(__func__, mul_node, /*num_src=*/2);
const ggml_tensor * x = unary_node->src[0];
const ggml_tensor * g = (mul_node->src[0] == unary_node) ? mul_node->src[1] : mul_node->src[0];
// g is picked by elimination; ggml_can_fuse()'s single-use rule rules out MUL(unary, unary)
GGML_ASSERT(g != unary_node);
GGML_ASSERT(x->type == g->type && x->type == mul_node->type);
GGML_ASSERT(ggml_are_same_shape(x, g) && ggml_are_same_shape(x, mul_node));
GGML_ASSERT(ggml_is_contiguous_1(x) && ggml_is_contiguous_1(g));
// dst is indexed flat
GGML_ASSERT(ggml_is_contiguous(mul_node));
queue_ptr main_stream = ctx.stream();
SYCL_CHECK(ggml_sycl_set_device(ctx.device));
const int64_t k = ggml_nelements(mul_node);
const int64_t n = mul_node->ne[0];
const auto dispatch_type = [&](auto op) {
switch (mul_node->type) {
case GGML_TYPE_F32:
unary_mul_sycl((const float *) x->data, (const float *) g->data, (float *) mul_node->data,
k, n, x->nb[1] / sizeof(float), g->nb[1] / sizeof(float), main_stream, op);
break;
case GGML_TYPE_F16:
unary_mul_sycl((const sycl::half *) x->data, (const sycl::half *) g->data, (sycl::half *) mul_node->data,
k, n, x->nb[1] / sizeof(sycl::half), g->nb[1] / sizeof(sycl::half), main_stream, op);
break;
default:
GGML_ABORT("fused unary+mul: unsupported type %s", ggml_type_name(mul_node->type));
}
};
switch (ggml_get_unary_op(unary_node)) {
case GGML_UNARY_OP_SILU: dispatch_type([](float v) { return op_silu(v); }); break;
case GGML_UNARY_OP_SIGMOID: dispatch_type([](float v) { return op_sigmoid(v); }); break;
case GGML_UNARY_OP_SOFTPLUS: dispatch_type([](float v) { return op_softplus(v); }); break;
default:
GGML_ABORT("fused unary+mul: unsupported unary op %s", ggml_unary_op_name(ggml_get_unary_op(unary_node)));
}
}
__dpct_inline__ float ggml_sycl_op_swiglu_oai_single(float x, float g, float alpha = 1.702f, float limit = 7.0f) {
x = sycl::fmin(x, limit);
g = sycl::fmax(sycl::fmin(g, limit), -limit);