Skip to content

sycl: *glu flat path #26354

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
ggerganov merged 3 commits into from
Aug 7, 2026
Merged

sycl: *glu flat path #26354

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
f242277
tests: add SWIGLU perf cases
Titaniumtown Jul 30, 2026
9357992
sycl: consolidate fused-GLU kernels
Titaniumtown Jul 30, 2026
154aec0
sycl: contiguous fast path for the fused GLU ops
Titaniumtown Jul 30, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
148 changes: 60 additions & 88 deletions ggml/src/ggml-sycl/element_wise.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -420,53 +420,31 @@ static void clamp(const T * x, T * dst, const float min, const float max, const
}
}

template<typename T>
static void gated_op_fused_geglu(const T * x, const T * g, T * dst, const uint64_t k, const sycl::uint3 n_fd, const uint64_t o0, const uint64_t o1, const sycl::nd_item<1> &item_ct1) {
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] = op_gelu(x[j0]) * g[j1];
}
}

template<typename T>
static void gated_op_fused_reglu(const T * x, const T * g, T * dst, const uint64_t k, const sycl::uint3 n_fd, const uint64_t o0, const uint64_t o1, const sycl::nd_item<1> &item_ct1) {
template<typename T, typename F>
static void unary_gated_op_flat_kernel(const T * x, const T * g, T * dst, const uint64_t k, const sycl::nd_item<1> & item_ct1, F func) {
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] = op_relu(x[j0]) * g[j1];
dst[i] = func(x[i]) * g[i];
}
}

template<typename T>
static void gated_op_fused_swiglu(const T * x, const T * g, T * dst, const uint64_t k, const sycl::uint3 n_fd, const uint64_t o0, const uint64_t o1, const sycl::nd_item<1> &item_ct1) {
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] = op_silu(x[j0]) * g[j1];
}
}

template<typename T>
static void gated_op_fused_geglu_erf(const T * x, const T * g, T * dst, const uint64_t k, const sycl::uint3 n_fd, const uint64_t o0, const uint64_t o1, const sycl::nd_item<1> &item_ct1) {
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] = op_gelu_erf(x[j0]) * g[j1];
}
}
template<typename T, typename F>
static void unary_gated_op_generic_kernel(
const T * x,
const T * g,
T * dst,
const uint64_t k,
const sycl::uint3 n_fd,
const uint64_t o0,
const uint64_t o1,
const sycl::nd_item<1> & item_ct1,
F func) {

template<typename T>
static void gated_op_fused_geglu_quick(const T * x, const T * g, T * dst, const uint64_t k, const sycl::uint3 n_fd, const uint64_t o0, const uint64_t o1, const sycl::nd_item<1> &item_ct1) {
// rows of n columns at strides o0 and o1: two halves of one fused tensor, or two tensors
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] = op_gelu_quick(x[j0]) * g[j1];
dst[i] = func(x[j0]) * g[j1];
}
}

Expand Down Expand Up @@ -670,6 +648,35 @@ static inline void ggml_sycl_op_unary(
});
}

template<typename F>
static inline void ggml_sycl_op_unary_gated(
ggml_backend_sycl_context & ctx, ggml_tensor * dst, F func) {

dispatch_ggml_sycl_op_fused_glu(ctx, dst,
[func](const auto * x_ptr, const auto * g_ptr, auto * dst_ptr, uint64_t k, uint64_t n, uint64_t o0, uint64_t o1, queue_ptr main_stream) {

const uint32_t num_blocks = (uint32_t) ceil_div(k, SYCL_GLU_BLOCK_SIZE);
const sycl::nd_range<1> launch_range(num_blocks * sycl::range<1>(SYCL_GLU_BLOCK_SIZE),
sycl::range<1>(SYCL_GLU_BLOCK_SIZE));

// o0 == n and o1 == n make the index math the identity, so index flat
// note: not ggml_is_contiguous - a fused [gate|up] src0 is contiguous with o0 == 2n
if (o0 == n && o1 == n) {
main_stream->parallel_for(launch_range,
[=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
unary_gated_op_flat_kernel(x_ptr, g_ptr, dst_ptr, k, item_ct1, func);
});
} else {
// launch-invariant divisor, and only this path needs it
const sycl::uint3 n_fd = init_fastdiv_values((uint32_t) n);
main_stream->parallel_for(launch_range,
[=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
unary_gated_op_generic_kernel(x_ptr, g_ptr, dst_ptr, k, n_fd, o0, o1, item_ct1, func);
});
}
});
}


static inline void ggml_sycl_op_arange(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
GGML_ASSERT(dst->type == GGML_TYPE_F32);
Expand Down Expand Up @@ -967,42 +974,21 @@ static inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, ggml_tensor
}

static inline void ggml_sycl_op_geglu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
ggml_sycl_detail::dispatch_ggml_sycl_op_fused_glu(ctx, dst,
[](const auto* x_ptr, const auto* g_ptr, auto* dst_ptr, uint64_t k, uint64_t n, uint64_t o0, uint64_t o1, queue_ptr main_stream) {
const uint32_t num_blocks = ceil_div(k, SYCL_GELU_BLOCK_SIZE);
const sycl::uint3 n_fd = init_fastdiv_values((uint32_t) n);
main_stream->parallel_for(
sycl::nd_range<1>((num_blocks * sycl::range<1>(SYCL_GELU_BLOCK_SIZE)),
sycl::range<1>(SYCL_GELU_BLOCK_SIZE)), [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
gated_op_fused_geglu(x_ptr, g_ptr, dst_ptr, k, n_fd, o0, o1, item_ct1);
});
});
ggml_sycl_detail::ggml_sycl_op_unary_gated(ctx, dst, [](auto x) {
return op_gelu(x);
});
}

static inline void ggml_sycl_op_reglu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
ggml_sycl_detail::dispatch_ggml_sycl_op_fused_glu(ctx, dst,
[](const auto* x_ptr, const auto* g_ptr, auto* dst_ptr, uint64_t k, uint64_t n, uint64_t o0, uint64_t o1, queue_ptr main_stream) {
const uint32_t num_blocks = ceil_div((uint32_t)k, SYCL_RELU_BLOCK_SIZE); // Using RELU block size for reglu
const sycl::uint3 n_fd = init_fastdiv_values((uint32_t) n);
main_stream->parallel_for(
sycl::nd_range<1>((num_blocks * sycl::range<1>(SYCL_RELU_BLOCK_SIZE)),
sycl::range<1>(SYCL_RELU_BLOCK_SIZE)), [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
gated_op_fused_reglu(x_ptr, g_ptr, dst_ptr, k, n_fd, o0, o1, item_ct1);
});
});
ggml_sycl_detail::ggml_sycl_op_unary_gated(ctx, dst, [](auto x) {
return op_relu(x);
});
}

static inline void ggml_sycl_op_swiglu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
ggml_sycl_detail::dispatch_ggml_sycl_op_fused_glu(ctx, dst,
[](const auto* x_ptr, const auto* g_ptr, auto* dst_ptr, uint64_t k, uint64_t n, uint64_t o0, uint64_t o1, queue_ptr main_stream) {
const uint32_t num_blocks = ceil_div((uint32_t)k, SYCL_SILU_BLOCK_SIZE); // Using SILU block size for swiglu
const sycl::uint3 n_fd = init_fastdiv_values((uint32_t) n);
main_stream->parallel_for(
sycl::nd_range<1>((num_blocks * sycl::range<1>(SYCL_SILU_BLOCK_SIZE)),
sycl::range<1>(SYCL_SILU_BLOCK_SIZE)), [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
gated_op_fused_swiglu(x_ptr, g_ptr, dst_ptr, k, n_fd, o0, o1, item_ct1);
});
});
ggml_sycl_detail::ggml_sycl_op_unary_gated(ctx, dst, [](auto x) {
return op_silu(x);
});
}

__dpct_inline__ float ggml_sycl_op_swiglu_oai_single(float x, float g, float alpha = 1.702f, float limit = 7.0f) {
Expand Down Expand Up @@ -1097,29 +1083,15 @@ void ggml_sycl_op_swiglu_oai(ggml_backend_sycl_context & ctx, ggml_tensor * dst)
}

static inline void ggml_sycl_op_geglu_erf(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
ggml_sycl_detail::dispatch_ggml_sycl_op_fused_glu(ctx, dst,
[](const auto* x_ptr, const auto* g_ptr, auto* dst_ptr, uint64_t k, uint64_t n, uint64_t o0, uint64_t o1, queue_ptr main_stream) {
const uint32_t num_blocks = ceil_div(k, SYCL_GELU_BLOCK_SIZE);
const sycl::uint3 n_fd = init_fastdiv_values((uint32_t) n);
main_stream->parallel_for(
sycl::nd_range<1>((num_blocks * sycl::range<1>(SYCL_GELU_BLOCK_SIZE)),
sycl::range<1>(SYCL_GELU_BLOCK_SIZE)), [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
gated_op_fused_geglu_erf(x_ptr, g_ptr, dst_ptr, k, n_fd, o0, o1, item_ct1);
});
});
ggml_sycl_detail::ggml_sycl_op_unary_gated(ctx, dst, [](auto x) {
return op_gelu_erf(x);
});
}

static inline void ggml_sycl_op_geglu_quick(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
ggml_sycl_detail::dispatch_ggml_sycl_op_fused_glu(ctx, dst,
[](const auto* x_ptr, const auto* g_ptr, auto* dst_ptr, uint64_t k, uint64_t n, uint64_t o0, uint64_t o1, queue_ptr main_stream) {
const uint32_t num_blocks = ceil_div(k, SYCL_GELU_BLOCK_SIZE);
const sycl::uint3 n_fd = init_fastdiv_values((uint32_t) n);
main_stream->parallel_for(
sycl::nd_range<1>((num_blocks * sycl::range<1>(SYCL_GELU_BLOCK_SIZE)),
sycl::range<1>(SYCL_GELU_BLOCK_SIZE)), [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
gated_op_fused_geglu_quick(x_ptr, g_ptr, dst_ptr, k, n_fd, o0, o1, item_ct1);
});
});
ggml_sycl_detail::ggml_sycl_op_unary_gated(ctx, dst, [](auto x) {
return op_gelu_quick(x);
});
}


Expand Down
2 changes: 0 additions & 2 deletions ggml/src/ggml-sycl/presets.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,8 +20,6 @@
#define MATRIX_ROW_PADDING 512 // last row of quant. matrices is a multiple of this to avoid out-of-bounds memory accesses

#define SYCL_COL2IM_1D_BLOCK_SIZE 256
#define SYCL_GELU_BLOCK_SIZE 256
#define SYCL_SILU_BLOCK_SIZE 256
#define SYCL_TANH_BLOCK_SIZE 256
#define SYCL_RELU_BLOCK_SIZE 256
#define SYCL_HARDSIGMOID_BLOCK_SIZE 256
Expand Down
9 changes: 9 additions & 0 deletions tests/test-backend-ops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9726,6 +9726,15 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
std::vector<std::unique_ptr<test_case>> test_cases;

// SWIGLU at a 27B-class FFN width, fused [gate|up] vs split operands
// note: same bytes either way, so a backend that indexes them differently shows it here
for (ggml_type type : {GGML_TYPE_F16, GGML_TYPE_F32}) {
for (int64_t n_tokens : {512, 2048}) {
test_cases.emplace_back(new test_glu(GGML_GLU_OP_SWIGLU, type, { 2*17408, n_tokens, 1, 1 }, 0, false));
test_cases.emplace_back(new test_glu_split(GGML_GLU_OP_SWIGLU, type, { 17408, n_tokens, 1, 1 }, 0));
}
}

// Conv2d: K=CRS=NPQ=4096 matmul performance
uint32_t iwh_idx = 0;
uint32_t kwh_idx = 1;
Expand Down
Loading