ikawrakow/ik_llama.cpp
1
0
Fork 0
mirror of https://github.com/ikawrakow/ik_llama.cpp.git synced 2026-02-23 22:54:10 +00:00
Code Issues Packages Projects Releases Wiki Activity

This works, but is slower than the non-working version

Browse Source
This commit is contained in:
Iwan Kawrakow
2025-07-05 20:09:15 +03:00
parent 030ba3aebf
commit b931e8b831
1 changed files with 13 additions and 35 deletions
Download Patch File Download Diff File Expand all files Collapse all files

48
ggml/src/ggml-cuda.cu
View File

@@ -2186,6 +2186,7 @@ struct mmid_row_mapping {
int32_t i2;
};
template <typename data_t = float>
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<char> 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<char> src1_contiguous(ctx.pool(), sizeof(float)*ggml_nelements(src1));
ggml_cuda_pool_alloc<char> dst_up_contiguous(ctx.pool(), sizeof(float)*ggml_nelements(dst));
ggml_cuda_pool_alloc<char> dst_gate_contiguous(ctx.pool(), sizeof(float)*ggml_nelements(dst));
ggml_cuda_pool_alloc<char> 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<<<grid_dims, block_dims, 0, stream>>>(
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<<<grid_dims, block_dims, 0, stream>>>(
@@ -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);