public_git/whisper.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/whisper.cpp.git synced 2026-10-22 12:10:55 +02:00
Code Issues Packages Projects Releases Wiki Activity

opencl: use buf+offset instead of subbuffer for adreno moe expert sort (llama/30418)

Browse Source
This commit is contained in:
lhez
2026-10-19 09:47:11 +02:00
committed by Daniel Bevenius
parent f562609708
commit 900f40e1a2
2 changed files with 18 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-opencl/ggml-opencl.cpp
+6 -8
View File
@@ -25925,10 +25925,6 @@ static void moe_router_reoerder(ggml_backend_t backend, const ggml_tensor * src,
const int max_post_router_tile = (ne20 * ne21 / n_tile_size) + ne02;
cl_buffer_region region;
region.origin = offset;
region.size = nb21 * ne21;
cl_mem original_router_buf = clCreateSubBuffer(extra->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, &region, &err);
CL_CHECK(err);
backend_ctx->prealloc_post_router.allocate(backend_ctx->context, sizeof(int) * max_post_router_tile * n_tile_size);
region.origin = 0;
@@ -25968,11 +25964,12 @@ static void moe_router_reoerder(ggml_backend_t backend, const ggml_tensor * src,
// Histogram
cl_kernel kernel = backend_ctx->kernel_moe_histogram;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &original_router_buf));
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra->data_device));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &hist_buf));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), &ne21));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), &ne20));
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), &ne02));
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &offset));
size_t histogram_global_size[] = {(size_t)(((ne21 + 63) / 64) * 64), static_cast<size_t>(ne20), 1};
size_t histogram_local_size[] = {64, 1, 1};
@@ -26012,13 +26009,14 @@ static void moe_router_reoerder(ggml_backend_t backend, const ggml_tensor * src,
if (stable_scatter) {
kernel = backend_ctx->kernel_moe_scatter_stable;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &original_router_buf));
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra->data_device));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &post_router_buf));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &emap_buf));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &tile_offset_buf));
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), &ne21));
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), &ne20));
CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &ne02));
CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_ulong), &offset));
// one workgroup (one wave) per expert; each ranks its own tokens
size_t scatter_global_size[] = {64, (size_t)ne02};
@@ -26026,7 +26024,7 @@ static void moe_router_reoerder(ggml_backend_t backend, const ggml_tensor * src,
backend_ctx->enqueue_ndrange_kernel(kernel, 2, scatter_global_size, scatter_local_size, src);
} else {
kernel = backend_ctx->kernel_moe_scatter;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &original_router_buf));
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra->data_device));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &post_router_buf));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &emap_buf));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &tile_offset_buf));
@@ -26034,6 +26032,7 @@ static void moe_router_reoerder(ggml_backend_t backend, const ggml_tensor * src,
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), &ne21));
CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &ne20));
CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &ne02));
CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_ulong), &offset));
backend_ctx->enqueue_ndrange_kernel(kernel, 3, histogram_global_size, histogram_local_size, src);
}
@@ -26055,7 +26054,6 @@ static void moe_router_reoerder(ggml_backend_t backend, const ggml_tensor * src,
fflush(stderr);
}
CL_CHECK(clReleaseMemObject(original_router_buf));
CL_CHECK(clReleaseMemObject(hist_buf));
CL_CHECK(clReleaseMemObject(tile_offset_buf));
CL_CHECK(clReleaseMemObject(total_tiles_buf));
ggml/src/ggml-opencl/kernels/moe_sort_by_expert.cl
+12 -3
View File
@@ -5,8 +5,11 @@ __kernel void kernel_moe_histogram(
__global int * hist,
uint N,
uint topK,
uint n_experts
uint n_experts,
ulong offset
) {
input = (__global const int *)((__global const char *)input + offset);
uint n = get_global_id(0);
uint k = get_global_id(1);
@@ -47,8 +50,11 @@ __kernel void kernel_moe_scatter(
__global int * slot_counter,
int N,
int topK,
uint n_experts
uint n_experts,
ulong offset
) {
input = (__global const int *)((__global const char *)input + offset);
uint n = get_global_id(0);
uint k = get_global_id(1);
@@ -89,8 +95,11 @@ __kernel void kernel_moe_scatter_stable(
__global const int * tile_offset,
int N,
int topK,
uint n_experts
uint n_experts,
ulong offset
) {
input = (__global const int *)((__global const char *)input + offset);
const int e = get_group_id(1);
const int lid = get_local_id(0);
const int M = N * topK;