From ac70ca35dd8928c10554ff3cd180974281924938 Mon Sep 17 00:00:00 2001 From: Concedo <39025047+LostRuins@users.noreply.github.com> Date: Sun, 22 Feb 2026 12:50:08 +0800 Subject: [PATCH] preliminary patches for acestep.cpp --- ggml/src/ggml-cuda/conv-transpose-1d.cu | 9 +++++---- ggml/src/ggml-cuda/im2col.cu | 5 +++-- 2 files changed, 8 insertions(+), 6 deletions(-) diff --git a/ggml/src/ggml-cuda/conv-transpose-1d.cu b/ggml/src/ggml-cuda/conv-transpose-1d.cu index 8418ba667..7dad163a5 100644 --- a/ggml/src/ggml-cuda/conv-transpose-1d.cu +++ b/ggml/src/ggml-cuda/conv-transpose-1d.cu @@ -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]; diff --git a/ggml/src/ggml-cuda/im2col.cu b/ggml/src/ggml-cuda/im2col.cu index 56dc05457..c1b223727 100644 --- a/ggml/src/ggml-cuda/im2col.cu +++ b/ggml/src/ggml-cuda/im2col.cu @@ -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<<>>(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);