diff --git a/ggml/src/ggml-sycl/element_wise.cpp b/ggml/src/ggml-sycl/element_wise.cpp index 0e707d531..11c94bcee 100644 --- a/ggml/src/ggml-sycl/element_wise.cpp +++ b/ggml/src/ggml-sycl/element_wise.cpp @@ -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 +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 +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 +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); diff --git a/ggml/src/ggml-sycl/element_wise.hpp b/ggml/src/ggml-sycl/element_wise.hpp index beea052cf..9f660f1b7 100644 --- a/ggml/src/ggml-sycl/element_wise.hpp +++ b/ggml/src/ggml-sycl/element_wise.hpp @@ -95,4 +95,7 @@ void ggml_sycl_trunc(ggml_backend_sycl_context & ctx, ggml_tensor * dst); void ggml_sycl_arange(ggml_backend_sycl_context & ctx, ggml_tensor * dst); +// fused UNARY(silu|sigmoid|softplus) + MUL; see ggml_sycl_can_fuse() for the accepted shapes +void ggml_sycl_op_unary_mul_fused(ggml_backend_sycl_context & ctx, ggml_tensor * unary_node, ggml_tensor * mul_node); + #endif // GGML_SYCL_ELEMENTWISE_HPP diff --git a/ggml/src/ggml-sycl/fusion.cpp b/ggml/src/ggml-sycl/fusion.cpp index 4a6027f39..97af2a1e4 100644 --- a/ggml/src/ggml-sycl/fusion.cpp +++ b/ggml/src/ggml-sycl/fusion.cpp @@ -1,6 +1,14 @@ #include "fusion.hpp" -bool ggml_sycl_can_fuse(const ggml_cgraph * cgraph, int node_idx, std::initializer_list ops) { +#include + +bool ggml_sycl_can_fuse(const ggml_cgraph * cgraph, int node_idx, std::initializer_list ops, + std::initializer_list unary_ops) { +#ifndef NDEBUG + const size_t num_unary = std::count(ops.begin(), ops.end(), GGML_OP_UNARY); + GGML_ASSERT(unary_ops.size() == num_unary); +#endif + if (!g_ggml_sycl_enable_fusion) { return false; } @@ -40,5 +48,45 @@ bool ggml_sycl_can_fuse(const ggml_cgraph * cgraph, int node_idx, std::initializ return true; } + if (ops.size() == 2 && ops.begin()[0] == GGML_OP_UNARY && ops.begin()[1] == GGML_OP_MUL && + unary_ops.size() == 1) { + const ggml_tensor * unary = cgraph->nodes[node_idx]; + const ggml_tensor * mul = cgraph->nodes[node_idx + 1]; + + const ggml_unary_op unary_op = ggml_get_unary_op(unary); + if (unary_op != unary_ops.begin()[0]) { + return false; + } + + // the ops ggml_sycl_op_unary_mul_fused() has a kernel for + if (unary_op != GGML_UNARY_OP_SILU && unary_op != GGML_UNARY_OP_SIGMOID && + unary_op != GGML_UNARY_OP_SOFTPLUS) { + return false; + } + + if (unary->type != GGML_TYPE_F32 && unary->type != GGML_TYPE_F16) { + return false; + } + + const ggml_tensor * other = (mul->src[0] == unary) ? mul->src[1] : mul->src[0]; + if (other->type != unary->type) { + return false; + } + + // one row stride per source comes from nb[1], so rows must be contiguous and equally + // shaped; the destination is written flat, so it must be fully contiguous + if (!ggml_is_contiguous_1(unary->src[0]) || !ggml_is_contiguous_1(other) || + !ggml_are_same_shape(other, unary) || !ggml_is_contiguous(mul)) { + return false; + } + + // the 32-bit fastdiv is inexact past 2^31; decline, the unfused path handles it + if (ggml_nelements(mul) >= ((int64_t) 1 << 31)) { + return false; + } + + return true; + } + return false; } diff --git a/ggml/src/ggml-sycl/fusion.hpp b/ggml/src/ggml-sycl/fusion.hpp index 7d7c79e02..94e74088c 100644 --- a/ggml/src/ggml-sycl/fusion.hpp +++ b/ggml/src/ggml-sycl/fusion.hpp @@ -6,10 +6,12 @@ #include "common.hpp" // Backend-side fusability test. `ops` names a candidate op sequence starting at cgraph node -// `node_idx`; the result is true only if ggml considers that subgraph fusable *and* the SYCL +// `node_idx`, and `unary_ops` the GGML_UNARY_OP each GGML_OP_UNARY in `ops` must carry, in +// order; the result is true only if ggml considers that subgraph fusable *and* the SYCL // kernel which would service it accepts the tensors involved (types, shapes, contiguity). // // Lives in its own translation unit because it grows a branch per supported op sequence. -bool ggml_sycl_can_fuse(const ggml_cgraph * cgraph, int node_idx, std::initializer_list ops); +bool ggml_sycl_can_fuse(const ggml_cgraph * cgraph, int node_idx, std::initializer_list ops, + std::initializer_list unary_ops); #endif // GGML_SYCL_FUSION_HPP diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index b98a28832..e712c18ba 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -5452,11 +5452,17 @@ static void ggml_backend_sycl_graph_compute_impl(ggml_backend_sycl_context * syc } #endif if (node->op == GGML_OP_RMS_NORM && - ggml_sycl_can_fuse(cgraph, i, { GGML_OP_RMS_NORM, GGML_OP_MUL })) { + ggml_sycl_can_fuse(cgraph, i, { GGML_OP_RMS_NORM, GGML_OP_MUL }, {})) { ggml_sycl_op_rms_norm_fused(*sycl_ctx, node, cgraph->nodes[i + 1]); i++; continue; } + if (node->op == GGML_OP_UNARY && + ggml_sycl_can_fuse(cgraph, i, { GGML_OP_UNARY, GGML_OP_MUL }, { ggml_get_unary_op(node) })) { + ggml_sycl_op_unary_mul_fused(*sycl_ctx, node, cgraph->nodes[i + 1]); + i++; + continue; + } bool ok = ggml_sycl_compute_forward(*sycl_ctx, node); if (!ok) { diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 631a0cfcd..693816542 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -3695,6 +3695,117 @@ struct test_relu_sqr : public test_case { } }; +// GGML_OP_UNARY(SILU|SIGMOID|SOFTPLUS) + GGML_OP_MUL (fused operation). +// `layout` and `tail` are used for fallback cases where fusion must be skipped +struct test_unary_mul : public test_case { + const ggml_unary_op op; + const ggml_type type; + const std::array ne; + const bool swap; // unary result is the second MUL operand + const std::string layout; // operand layout, see build_graph() + const std::string tail; // extra consumer past the MUL, see build_graph() + + std::string op_desc(ggml_tensor * t) override { + GGML_UNUSED(t); + return std::string(ggml_unary_op_name(op)) + "_MUL"; + } + + bool run_whole_graph() override { return true; } + + double max_nmse_err() override { + // the fused kernel elides the rounding of the unary result that the CPU chain + // performs; relax the tolerance to match that drift + switch (type) { + case GGML_TYPE_F16: return 5e-5; + default: return 1e-7; + } + } + + std::string vars() override { + return VARS_TO_STR5(type, ne, swap, layout, tail); + } + + test_unary_mul(ggml_unary_op op, + ggml_type type = GGML_TYPE_F32, + std::array ne = {128, 2, 2, 2}, + bool swap = false, + std::string layout = "packed", + std::string tail = "") + : op(op), type(type), ne(ne), swap(swap), layout(std::move(layout)), tail(std::move(tail)) {} + + // `ne` viewed out of a wider tensor: rows stay contiguous, but the stride exceeds the width + ggml_tensor * padded(ggml_context * ctx, const char * name, int64_t mul0, int64_t off0) { + std::array ne_w = ne; + ne_w[0] *= mul0; + ggml_tensor * base = ggml_new_tensor(ctx, type, 4, ne_w.data()); + ggml_set_name(base, name); + return ggml_view_4d(ctx, base, ne[0], ne[1], ne[2], ne[3], + base->nb[1], base->nb[2], base->nb[3], off0 * base->nb[0]); + } + + ggml_tensor * build_graph(ggml_context * ctx) override { + ggml_tensor * a = nullptr; // unary source + ggml_tensor * b = nullptr; // other MUL operand + + if (layout == "packed") { + a = ggml_new_tensor(ctx, type, 4, ne.data()); + b = ggml_new_tensor(ctx, type, 4, ne.data()); + } else if (layout == "pad_unary") { + a = padded(ctx, "a", 3, 0); + b = ggml_new_tensor(ctx, type, 4, ne.data()); + } else if (layout == "pad_other") { + a = ggml_new_tensor(ctx, type, 4, ne.data()); + b = padded(ctx, "b", 3, 0); + } else if (layout == "halves") { + // the shape the Conformer audio encoders build: one tensor split in two + std::array ne_w = ne; + ne_w[0] *= 2; + ggml_tensor * base = ggml_new_tensor(ctx, type, 4, ne_w.data()); + ggml_set_name(base, "base"); + b = ggml_view_4d(ctx, base, ne[0], ne[1], ne[2], ne[3], base->nb[1], base->nb[2], base->nb[3], 0); + a = ggml_view_4d(ctx, base, ne[0], ne[1], ne[2], ne[3], base->nb[1], base->nb[2], base->nb[3], + ne[0] * base->nb[0]); + } else if (layout == "strided_dim1") { + // contiguous rows but a strided dim 1: not ggml_is_contiguous_1, must not fuse + std::array ne_w = ne; + ne_w[1] *= 3; + ggml_tensor * base = ggml_new_tensor(ctx, type, 4, ne_w.data()); + ggml_set_name(base, "a"); + a = ggml_view_4d(ctx, base, ne[0], ne[1], ne[2], ne[3], base->nb[1], base->nb[2], base->nb[3], 0); + b = ggml_new_tensor(ctx, type, 4, ne.data()); + } else if (layout == "bcast") { + a = ggml_new_tensor(ctx, type, 4, ne.data()); + b = ggml_new_tensor_4d(ctx, type, ne[0], 1, 1, 1); + } else { + GGML_ABORT("unknown layout %s", layout.c_str()); + } + ggml_set_name(a, "a"); + ggml_set_name(b, "b"); + + ggml_tensor * u = ggml_unary(ctx, a, op); + ggml_set_name(u, "unary"); + + // a broadcasting operand can only be the second one + const bool second = swap && layout != "bcast"; + ggml_tensor * out = second ? ggml_mul(ctx, b, u) : ggml_mul(ctx, u, b); + + if (tail == "reuse") { + // a second read of the unary result must block the fusion + ggml_set_name(out, "mul"); + out = ggml_add(ctx, out, u); + } else if (tail == "consumer") { + // fusion still applies; catches a dispatcher that skips one node too many + ggml_set_name(out, "mul"); + out = ggml_add(ctx, out, b); + } else if (!tail.empty()) { + GGML_ABORT("unknown tail %s", tail.c_str()); + } + ggml_set_name(out, "out"); + + return out; + } +}; + // SNAKE activation fusion: y = x + sin(a*x)^2 * inv_b // CUDA backend matches the naive 5-op chain (mul, sin, sqr, mul, add) // and dispatches a single fused kernel. @@ -8065,6 +8176,25 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_relu_sqr(type, { 5, 7, 11, 13 })); } + // fused unary + mul (gated activations that are not expressed as GGML_OP_GLU) + for (ggml_unary_op op : { GGML_UNARY_OP_SILU, GGML_UNARY_OP_SIGMOID, GGML_UNARY_OP_SOFTPLUS }) { + for (ggml_type type : { GGML_TYPE_F16, GGML_TYPE_F32 }) { + for (bool swap : { false, true }) { + test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, swap)); + } + test_cases.emplace_back(new test_unary_mul(op, type, { 5, 7, 11, 13 })); + test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, false, "pad_unary")); + // a view only stays out from between the two ops when the unary result is second + test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, true, "pad_other")); + test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, true, "halves")); + test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, false, "packed", "consumer")); + // must not fuse + test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, false, "strided_dim1")); + test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, false, "bcast")); + test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, false, "packed", "reuse")); + } + } + // SNAKE activation fusion: x + sin(a*x)^2 * inv_b for (ggml_type type : { GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_BF16 }) { test_cases.emplace_back(new test_snake_fuse(type, { 5, 7, 1, 1})); // primes sub-block