Skip to content

[pull] master from ggml-org:master #1177

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
pull merged 5 commits into from
May 13, 2026
Merged

[pull] master from ggml-org:master #1177

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
527045b
flush the gpu profile timestamp before the queryset is overflowed (#2…
yomaytk May 13, 2026
1e4579f
opencl: fix crash when warming up MoE on Adreno (#22876)
lhez May 13, 2026
95d469a
server, webui: accept continue_final_message flag for vLLM API compat…
ServeurpersoCom May 13, 2026
ec562eb
opencl: add q5_0 and q5_1 MoE for Adreno (#22985)
shaofeiqi May 13, 2026
7f3f843
Fix for issue #22974. Cast intermediate results to float before addin…
scutler-nv May 13, 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
11 changes: 7 additions & 4 deletions ggml/src/ggml-cuda/allreduce.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -184,13 +184,15 @@ static __global__ void ggml_cuda_ar_kernel(
#pragma unroll
for (int k = 0; k < ELEMS_PER_VEC; ++k) {
const T_wire d_low = ggml_cuda_cast<T_wire>(sendbuf[off + k]);
recvbuf[off + k] = ggml_cuda_cast<T_dst>(d_low) + ggml_cuda_cast<T_dst>(wire[k]);
recvbuf[off + k] = ggml_cuda_cast<T_dst>(
ggml_cuda_cast<float>(d_low) + ggml_cuda_cast<float>(wire[k]));
}
}
if (bid == 0 && tid < count - tail) {
const T_wire d_low = ggml_cuda_cast<T_wire>(sendbuf[tail + tid]);
recvbuf[tail + tid] =
ggml_cuda_cast<T_dst>(d_low) + ggml_cuda_cast<T_dst>(host_other[tail + tid]);
recvbuf[tail + tid] = ggml_cuda_cast<T_dst>(
ggml_cuda_cast<float>(d_low) +
ggml_cuda_cast<float>(host_other[tail + tid]));
}
}
}
Expand All @@ -210,7 +212,8 @@ static __global__ void ggml_cuda_ar_add_kernel(
const int nt = gridDim.x * blockDim.x;
for (int i = tid; i < count; i += nt) {
const T_src d_low = ggml_cuda_cast<T_src>(dst[i]);
dst[i] = ggml_cuda_cast<T_dst>(d_low) + ggml_cuda_cast<T_dst>(src[i]);
dst[i] = ggml_cuda_cast<T_dst>(
ggml_cuda_cast<float>(d_low) + ggml_cuda_cast<float>(src[i]));
}
}

Expand Down
4 changes: 4 additions & 0 deletions ggml/src/ggml-opencl/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -106,6 +106,10 @@ set(GGML_OPENCL_KERNELS
gemv_moe_q4_0_f32_ns
gemm_moe_q4_1_f32_ns
gemv_moe_q4_1_f32_ns
gemm_moe_q5_0_f32_ns
gemv_moe_q5_0_f32_ns
gemm_moe_q5_1_f32_ns
gemv_moe_q5_1_f32_ns
gemm_moe_mxfp4_f32
gemv_moe_mxfp4_f32
gemm_moe_mxfp4_f32_ns
Expand Down
Loading
Loading