From b931e8b831614ceaa55fe7ff2af9c7605bf52baa Mon Sep 17 00:00:00 2001 From: Iwan Kawrakow Date: Sat, 5 Jul 2025 20:09:15 +0300 Subject: [PATCH] This works, but is slower than the non-working version --- ggml/src/ggml-cuda.cu | 48 ++++++++++++------------------------------- 1 file changed, 13 insertions(+), 35 deletions(-) diff --git a/ggml/src/ggml-cuda.cu b/ggml/src/ggml-cuda.cu index 2dc48d4d..dfe003a4 100644 --- a/ggml/src/ggml-cuda.cu +++ b/ggml/src/ggml-cuda.cu @@ -2186,6 +2186,7 @@ struct mmid_row_mapping { int32_t i2; }; +template static __global__ void k_copy_src_to_contiguous(const char * __restrict__ src_original, char * __restrict__ src_contiguous, const mmid_row_mapping * __restrict__ row_mapping, int64_t ne10, int64_t ne11, size_t nb11, size_t nb12) { @@ -2194,8 +2195,8 @@ static __global__ void k_copy_src_to_contiguous(const char * __restrict__ src_or const int32_t i11 = row_mapping[i].i1 % ne11; const int32_t i12 = row_mapping[i].i2; - float * src_row_contiguous = (float *)(src_contiguous + i*nb11); - const float * src_row_original = (const float *)(src_original + i11*nb11 + i12*nb12); + data_t * src_row_contiguous = (data_t *)(src_contiguous + i*nb11); + const data_t * src_row_original = (const data_t *)(src_original + i11*nb11 + i12*nb12); for (int j = threadIdx.x; j < ne10; j += blockDim.x) { src_row_contiguous[j] = src_row_original[j]; @@ -2682,16 +2683,9 @@ static bool ggml_cuda_up_gate_unary(ggml_backend_cuda_context & ctx, ggml_tensor src1_padded_row_size = src1_padded_num_cols/ggml_blck_size(GGML_TYPE_Q8_1)*ggml_type_size(GGML_TYPE_Q8_1); src1_quantized_size = src1_padded_row_size*src1->ne[2] + get_mmq_x_max_host(ggml_cuda_info().devices[ctx.device].cc)*sizeof(block_q8_1_mmq); src1_quantized.alloc(src1_quantized_size); - quantize_mmq_q8_1_cuda((const float *)src1->data, src1_quantized.get(), src1->ne[0], src1->ne[2], src1->ne[3], src1_padded_num_cols, src0_1->type, stream); - CUDA_CHECK(cudaGetLastError()); use_quantized_src1 = true; } - ggml_cuda_pool_alloc src1_contiguous(ctx.pool()); - if (use_quantized_src1) { - src1_contiguous.alloc(src1_quantized_size); - } else { - src1_contiguous.alloc(sizeof(float)*ggml_nelements(src1)); - } + ggml_cuda_pool_alloc src1_contiguous(ctx.pool(), sizeof(float)*ggml_nelements(src1)); ggml_cuda_pool_alloc dst_up_contiguous(ctx.pool(), sizeof(float)*ggml_nelements(dst)); ggml_cuda_pool_alloc dst_gate_contiguous(ctx.pool(), sizeof(float)*ggml_nelements(dst)); ggml_cuda_pool_alloc final_dst_contiguous(ctx.pool()); @@ -2722,17 +2716,7 @@ static bool ggml_cuda_up_gate_unary(ggml_backend_cuda_context & ctx, ggml_tensor if (num_src1_rows == 0) continue; size_t mapping_offset = cum_moe_counts[i02]; - if (use_quantized_src1) { - unsigned int eff_ne10 = src1_padded_row_size/sizeof(float); - dim3 block_dims(std::min(eff_ne10, 768u)); - dim3 grid_dims(num_src1_rows); - k_copy_src_to_contiguous<<>>( - src1_quantized.get(), src1_contiguous.get(), dev_row_mapping.get() + mapping_offset, eff_ne10, ne11, src1_padded_row_size, src1_padded_row_size); - CUDA_CHECK(cudaGetLastError()); - src1_row.nb[0] = sizeof(block_q8_1); - src1_row.type = GGML_TYPE_Q8_1; - } - else { + { dim3 block_dims(std::min((unsigned int)ne10, 768u)); dim3 grid_dims(num_src1_rows); k_copy_src_to_contiguous<<>>( @@ -2756,22 +2740,16 @@ static bool ggml_cuda_up_gate_unary(ggml_backend_cuda_context & ctx, ggml_tensor dst_row.nb[2] = num_src1_rows*nb1; dst_row.nb[3] = num_src1_rows*nb1; -//struct mmq_args { -// const char * x; const char * y; float * dst; -// int64_t ne00; int64_t ne01; int64_t stride01; -// int64_t ne10; int64_t ne11; int64_t stride11; -// int64_t ne0; -//}; - -// const mmq_args args = {src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, nb01, src1_padded_row_size, src1_ncols, ne11, nrows_dst}; - - //ggml_cuda_op_mul_mat_vec_q_id(ctx, src0_1, &local_src1, ids, &local_dst, - // (const char *)src0_1->data, (const float *)src1->data, src1_quantized.get(), (float *)dst_up_contiguous.get(), - // 0, src0_1->ne[1], 1, src1_padded_col_size, stream); + if (use_quantized_src1) { + quantize_mmq_q8_1_cuda((const float *)src1_contiguous.get(), src1_quantized.get(), src1->ne[0], num_src1_rows, 1, + src1_padded_num_cols, src0_1->type, stream); + CUDA_CHECK(cudaGetLastError()); + src1_row.data = src1_quantized.get(); + } dst_row.data = dst_up_contiguous.get(); if (use_quantized_src1) { - ggml_cuda_op_mul_mat_q(ctx, &src0_1_row, &src1_row, &dst_row, (const char *)src0_1_row.data, nullptr, src1_contiguous.get(), (float *)dst_row.data, + ggml_cuda_op_mul_mat_q(ctx, &src0_1_row, &src1_row, &dst_row, (const char *)src0_1_row.data, nullptr, src1_quantized.get(), (float *)dst_row.data, 0, src0_1_row.ne[1], num_src1_rows, src1_padded_num_cols, stream); } else { ggml_cuda_mul_mat(ctx, &src0_1_row, &src1_row, &dst_row); @@ -2780,7 +2758,7 @@ static bool ggml_cuda_up_gate_unary(ggml_backend_cuda_context & ctx, ggml_tensor dst_row.data = dst_gate_contiguous.get(); if (use_quantized_src1) { - ggml_cuda_op_mul_mat_q(ctx, &src0_2_row, &src1_row, &dst_row, (const char *)src0_2_row.data, nullptr, src1_contiguous.get(), (float *)dst_row.data, + ggml_cuda_op_mul_mat_q(ctx, &src0_2_row, &src1_row, &dst_row, (const char *)src0_2_row.data, nullptr, src1_quantized.get(), (float *)dst_row.data, 0, src0_2_row.ne[1], num_src1_rows, src1_padded_num_cols, stream); } else { ggml_cuda_mul_mat(ctx, &src0_2_row, &src1_row, &dst_row);