vrr/koboldcpp
2
0
Fork 0
mirror of https://github.com/LostRuins/koboldcpp.git synced 2026-05-16 19:59:16 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions 3

revert to use upstream fix

Browse source
This commit is contained in:
Concedo 2026-05-14 18:56:57 +08:00
parent 4cfa1ad1c4
commit 6e42e0ebb6
1 changed files with 4 additions and 14 deletions
Download patch file Download diff file Expand all files Collapse all files

18
ggml/src/ggml-cuda/allreduce.cu
View file

@ -184,21 +184,13 @@ 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]);
const float a = ggml_cuda_cast<float>(d_low);
const float b = ggml_cuda_cast<float>(wire[k]);
recvbuf[off + k] = ggml_cuda_cast<T_dst>(a + b);
recvbuf[off + k] = ggml_cuda_cast<T_dst>(d_low) + ggml_cuda_cast<T_dst>(wire[k]);
}
}
if (bid == 0 && tid < count - tail) {
const T_wire d_low = ggml_cuda_cast<T_wire>(sendbuf[tail + tid]);
const float a = ggml_cuda_cast<float>(d_low);
const float b = ggml_cuda_cast<float>(host_other[tail + tid]);
recvbuf[tail + tid] = ggml_cuda_cast<T_dst>(a + b);
recvbuf[tail + tid] =
ggml_cuda_cast<T_dst>(d_low) + ggml_cuda_cast<T_dst>(host_other[tail + tid]);
}
}
}
@ -218,9 +210,7 @@ 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]);
const float a = (float) ggml_cuda_cast<float>(d_low);
const float b = (float) ggml_cuda_cast<float>(src[i]);
dst[i] = ggml_cuda_cast<T_dst>(a + b);
dst[i] = ggml_cuda_cast<T_dst>(d_low) + ggml_cuda_cast<T_dst>(src[i]);
}
}