vrr/koboldcpp
2
0
Fork 0
mirror of https://github.com/LostRuins/koboldcpp.git synced 2026-05-20 01:09:57 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

preliminary patches for acestep.cpp

Browse source
This commit is contained in:
Concedo 2026-02-22 12:50:08 +08:00
parent 19588f18ea
commit ac70ca35dd
2 changed files with 8 additions and 6 deletions
Download patch file Download diff file Expand all files Collapse all files

9
ggml/src/ggml-cuda/conv-transpose-1d.cu
View file

@ -21,10 +21,11 @@ static __global__ void conv_transpose_1d_kernel(
int kernel_offset = (src0_ne0 * src0_ne1 * c) + (out_index * src0_ne0);
int input_offset = src1_ne0 * c;
for (int i = 0; i < src1_ne0; i++) {
if (!(idx >= i*s0 && idx < i*s0 + src0_ne0)) {
continue;
}
int i_min = (idx >= src0_ne0) ? ((idx - src0_ne0 + s0) / s0) : 0;
int i_max_val = idx / s0;
int i_max = (i_max_val < src1_ne0) ? i_max_val : (src1_ne0 - 1);
for (int i = i_min; i <= i_max; i++) {
int weight_idx = idx - i*s0;
float kernel_weight = src0[kernel_offset + weight_idx];

5
ggml/src/ggml-cuda/im2col.cu
View file

@ -18,7 +18,7 @@ static __global__ void im2col_kernel(
const int64_t ikh = rem / KW;
const int64_t ikw = rem - ikh * KW;
const int64_t iow = blockIdx.y;
for (int64_t iow = blockIdx.y; iow < OW; iow += gridDim.y) {
for (int64_t iz = blockIdx.z; iz < N_OH; iz+=MAX_GRIDDIM_Z) {
const int64_t in = iz / OH;
const int64_t ioh = iz - in * OH;
@ -36,6 +36,7 @@ static __global__ void im2col_kernel(
dst[offset_dst] = x[offset_src + iih * IW + iiw];
}
}
}
GGML_UNUSED(IC);
GGML_UNUSED(KH);
@ -51,7 +52,7 @@ static void im2col_cuda(const float * x, T* dst,
const int64_t num_blocks = (IC_KH_KW + CUDA_IM2COL_BLOCK_SIZE - 1) / CUDA_IM2COL_BLOCK_SIZE;
const int64_t N_OH = N * OH;
const int64_t KH_KW = KW*KH;
dim3 block_nums(num_blocks, OW, MIN(N_OH, MAX_GRIDDIM_Z));
dim3 block_nums(num_blocks, MIN(OW, MAX_GRIDDIM_Z), MIN(N_OH, MAX_GRIDDIM_Z));
im2col_kernel<<<block_nums, MIN(IC_KH_KW, CUDA_IM2COL_BLOCK_SIZE) , 0, stream>>>(x, dst, IC, IW, IH, OH, OW, KW, KH,
IC_IH_IW, IH_IW, N_OH, KH_KW, IC_KH_KW,
s0, s1, p0, p1, d0, d1);