Skip to content

Commit 8d72972

Browse files
NexesenexNexesenex
committed
Revert "Even more fused ops (ikawrakow#868)"
This reverts commit bdf4f0d.
1 parent f59c391 commit 8d72972

File tree

6 files changed

+15
-159
lines changed

6 files changed

+15
-159
lines changed

‎ggml/src/ggml-cuda.cu‎

Lines changed: 5 additions & 52 deletions
Original file line numberDiff line numberDiff line change
@@ -2078,43 +2078,9 @@ static int ggml_cuda_mul_mat_q(ggml_backend_cuda_context & ctx, const ggml_tenso
20782078
src0->type, stream);
20792079
CUDA_CHECK(cudaGetLastError());
20802080

2081-
// The code below handles the case when Q, K, V have a bias applied after the resepctive matrix multiplication.
2082-
// In that case the graph contains mul_mat(Q) -> mul_mat(K) -> mul_mat(V) -> add(Q) -> add(K) -> add(V)
2083-
if (cgraph && node_n + 5 < cgraph->n_nodes &&
2084-
cgraph->nodes[node_n+1]->op == GGML_OP_MUL_MAT &&
2085-
cgraph->nodes[node_n+2]->op == GGML_OP_MUL_MAT &&
2086-
ggml_is_quantized(cgraph->nodes[node_n+1]->src[0]->type) &&
2087-
ggml_is_quantized(cgraph->nodes[node_n+2]->src[0]->type) &&
2088-
cgraph->nodes[node_n+3]->op == GGML_OP_ADD &&
2089-
cgraph->nodes[node_n+4]->op == GGML_OP_ADD &&
2090-
cgraph->nodes[node_n+5]->op == GGML_OP_ADD &&
2091-
cgraph->nodes[node_n+0] == cgraph->nodes[node_n+3]->src[0] &&
2092-
cgraph->nodes[node_n+1] == cgraph->nodes[node_n+4]->src[0] &&
2093-
cgraph->nodes[node_n+2] == cgraph->nodes[node_n+5]->src[0]) {
2094-
for (int i = 0; i < 3; ++i) {
2095-
auto src0_i = cgraph->nodes[node_n+i]->src[0];
2096-
ggml_cuda_op_mul_mat_vec_q_biased(ctx, src0_i, src1, cgraph->nodes[node_n+i], cgraph->nodes[node_n+i+3]->src[1],
2097-
(const char *)src0_i->data, nullptr, src1_quantized.get(), (float *)cgraph->nodes[node_n+i]->data,
2098-
0, src0_i->ne[1], src1->ne[1], ne10_padded, stream);
2099-
CUDA_CHECK(cudaGetLastError());
2100-
}
2101-
node_n += 5;
2102-
} else if (cgraph && node_n + 1 < cgraph->n_nodes &&
2103-
cgraph->nodes[node_n+1]->op == GGML_OP_ADD &&
2104-
dst == cgraph->nodes[node_n+1]->src[0] &&
2105-
dst->ne[0] == cgraph->nodes[node_n+1]->src[1]->ne[0] &&
2106-
cgraph->nodes[node_n+1]->src[1]->type == GGML_TYPE_F32 &&
2107-
ggml_nrows(cgraph->nodes[node_n+1]->src[1]) == 1) {
2108-
// We have a bias applied after the matrix multiplication and we can fuse it
2109-
ggml_cuda_op_mul_mat_vec_q_biased(ctx, dst->src[0], src1, cgraph->nodes[node_n+1], cgraph->nodes[node_n+1]->src[1],
2110-
(const char *)dst->src[0]->data, nullptr, src1_quantized.get(), (float *)cgraph->nodes[node_n+1]->data,
2111-
0, dst->src[0]->ne[1], src1->ne[1], ne10_padded, stream);
2112-
++node_n;
2113-
} else {
2114-
ggml_cuda_op_mul_mat_vec_q(ctx, src0, src1, dst, (const char *)src0->data, nullptr, src1_quantized.get(), (float *)dst->data,
2115-
0, src0->ne[1], src1->ne[1], ne10_padded, stream);
2116-
CUDA_CHECK(cudaGetLastError());
2117-
}
2081+
ggml_cuda_op_mul_mat_vec_q(ctx, src0, src1, dst, (const char *)src0->data, nullptr, src1_quantized.get(), (float *)dst->data,
2082+
0, src0->ne[1], src1->ne[1], ne10_padded, stream);
2083+
CUDA_CHECK(cudaGetLastError());
21182084
} else {
21192085
quantize_mmq_q8_1_cuda((const float *)src1->data, src1_quantized.get(), src1->ne[0], src1->ne[1], 1, ne10_padded, src0->type, stream);
21202086
CUDA_CHECK(cudaGetLastError());
@@ -2135,21 +2101,8 @@ static int ggml_cuda_mul_mat_q(ggml_backend_cuda_context & ctx, const ggml_tenso
21352101
if (dst->op != GGML_OP_MUL_MAT || dst->src[1] != src1 || !ggml_is_quantized(dst->src[0]->type)) break;
21362102
if (!is_gemv && mmq_get_q8_1_ds_layout(src0->type) != mmq_get_q8_1_ds_layout(dst->src[0]->type)) break;
21372103
if (is_gemv) {
2138-
if (node_n + 1 < cgraph->n_nodes &&
2139-
cgraph->nodes[node_n+1]->op == GGML_OP_ADD &&
2140-
dst == cgraph->nodes[node_n+1]->src[0] &&
2141-
dst->ne[0] == cgraph->nodes[node_n+1]->src[1]->ne[0] &&
2142-
cgraph->nodes[node_n+1]->src[1]->type == GGML_TYPE_F32 &&
2143-
ggml_nrows(cgraph->nodes[node_n+1]->src[1]) == 1) {
2144-
// We have a bias applied after the matrix multiplication and we can fuse it
2145-
ggml_cuda_op_mul_mat_vec_q_biased(ctx, dst->src[0], src1, cgraph->nodes[node_n+1], cgraph->nodes[node_n+1]->src[1],
2146-
(const char *)dst->src[0]->data, nullptr, src1_quantized.get(), (float *)cgraph->nodes[node_n+1]->data,
2147-
0, dst->src[0]->ne[1], src1->ne[1], ne10_padded, stream);
2148-
++node_n;
2149-
} else {
2150-
ggml_cuda_op_mul_mat_vec_q(ctx, dst->src[0], src1, dst, (const char *)dst->src[0]->data, nullptr, src1_quantized.get(),
2151-
(float *)dst->data, 0, dst->src[0]->ne[1], src1->ne[1], ne10_padded, stream);
2152-
}
2104+
ggml_cuda_op_mul_mat_vec_q(ctx, dst->src[0], src1, dst, (const char *)dst->src[0]->data, nullptr, src1_quantized.get(),
2105+
(float *)dst->data, 0, dst->src[0]->ne[1], src1->ne[1], ne10_padded, stream);
21532106
} else {
21542107
ggml_cuda_op_mul_mat_q(ctx, dst->src[0], src1, dst, (const char *)dst->src[0]->data, nullptr, src1_quantized.get(),
21552108
(float *)dst->data, 0, dst->src[0]->ne[1], src1->ne[1], ne10_padded, stream);

‎ggml/src/ggml-cuda/binbcast.cu‎

Lines changed: 0 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -313,25 +313,7 @@ void ggml_cuda_op_repeat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
313313
ggml_cuda_op_bin_bcast<bin_bcast_cuda<op_repeat>>(&aux_dst, &aux_src, &aux_dst, nullptr, dst->src[0]->data, dst->data, ctx.stream());
314314
}
315315

316-
static __global__ void k_fast_add(int64_t ne0, int64_t nelem, const float * x, const float * y, float * z) {
317-
int64_t i = blockDim.x*blockIdx.x + threadIdx.x;
318-
if (i >= nelem) {
319-
return;
320-
}
321-
z[i] = x[i] + y[i % ne0];
322-
}
323-
324316
void ggml_cuda_op_add(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
325-
if (ggml_nrows(dst->src[1]) == 1 && dst->src[0]->ne[0] == dst->src[1]->ne[0] &&
326-
dst->type == GGML_TYPE_F32 && dst->src[0]->type == GGML_TYPE_F32 && dst->src[1]->type == GGML_TYPE_F32 &&
327-
ggml_are_same_shape(dst, dst->src[0]) && ggml_is_contiguous(dst)) {
328-
constexpr int kBlockSize = 256;
329-
auto nelem = ggml_nelements(dst);
330-
int nblocks = (nelem + kBlockSize - 1)/kBlockSize;
331-
k_fast_add<<<nblocks, kBlockSize, 0, ctx.stream()>>>(dst->ne[0], nelem,
332-
(const float *)dst->src[0]->data, (const float *)dst->src[1]->data, (float *)dst->data);
333-
return;
334-
}
335317
ggml_cuda_op_bin_bcast<bin_bcast_cuda<op_add>>(dst->src[0], dst->src[1], dst, dst->src[0]->data, dst->src[1]->data, dst->data, ctx.stream());
336318
}
337319

‎ggml/src/ggml-cuda/cpy.cu‎

Lines changed: 6 additions & 50 deletions
Original file line numberDiff line numberDiff line change
@@ -38,25 +38,6 @@ static __global__ void cpy_flt(const char * cx, char * cdst_direct, const int ne
3838
cpy_1(cx + x_offset, cdst + dst_offset);
3939
}
4040

41-
template <typename src_t, typename dst_t>
42-
static __global__ void cpy_flt_contiguous(const char * cx, char * cdst_direct, const int ne,
43-
char ** cdst_indirect, int graph_cpynode_index) {
44-
const int64_t i = blockDim.x*blockIdx.x + threadIdx.x;
45-
46-
if (i >= ne) {
47-
return;
48-
}
49-
50-
auto dst = (cdst_indirect != nullptr) ? (dst_t *)cdst_indirect[graph_cpynode_index] : (dst_t *)cdst_direct;
51-
auto src = (const src_t *)cx;
52-
53-
if constexpr (std::is_same_v<dst_t, nv_bfloat16>) {
54-
dst[i] = __float2bfloat16(src[i]);
55-
} else {
56-
dst[i] = (dst_t)src[i];
57-
}
58-
}
59-
6041
static __device__ void cpy_blck_q8_0_f32(const char * cxi, char * cdsti) {
6142
float * cdstf = (float *)(cdsti);
6243

@@ -182,16 +163,6 @@ static void ggml_cpy_flt_cuda(
182163
(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, cdst_indirect, graph_cpynode_index++);
183164
}
184165

185-
template<typename src_t, typename dst_t>
186-
static void ggml_cpy_flt_contiguous_cuda(
187-
const char * cx, char * cdst, const int ne,
188-
cudaStream_t stream, char ** cdst_indirect, int & graph_cpynode_index) {
189-
190-
const int num_blocks = (ne + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE;
191-
cpy_flt_contiguous<src_t, dst_t><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>
192-
(cx, cdst, ne, cdst_indirect, graph_cpynode_index++);
193-
}
194-
195166
static void ggml_cpy_f32_q8_0_cuda(
196167
const char * cx, char * cdst, const int ne,
197168
const int ne00, const int ne01, const int ne02, const int nb00, const int nb01, const int nb02,
@@ -433,8 +404,6 @@ void ggml_cuda_cpy(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, gg
433404
char * src0_ddc = (char *) src0->data;
434405
char * src1_ddc = (char *) src1->data;
435406

436-
bool fast_cpy = ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && ggml_are_same_shape(src0, src1);
437-
438407
char ** dest_ptrs_d = nullptr;
439408
int graph_cpynode_index = -1;
440409
#if defined(GGML_CUDA_USE_GRAPHS) || defined(GGML_HIP_GRAPHS) || defined(GGML_MUSA_GRAPHS)
@@ -460,23 +429,11 @@ void ggml_cuda_cpy(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, gg
460429
}
461430
}
462431
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32) {
463-
if (fast_cpy) {
464-
ggml_cpy_flt_contiguous_cuda<float, float>(src0_ddc, src1_ddc, ne, main_stream, dest_ptrs_d, graph_cpynode_index);
465-
} else {
466-
ggml_cpy_flt_cuda<float, float> (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream, dest_ptrs_d, graph_cpynode_index);
467-
}
432+
ggml_cpy_flt_cuda<float, float> (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream, dest_ptrs_d, graph_cpynode_index);
468433
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_BF16) {
469-
if (fast_cpy) {
470-
ggml_cpy_flt_contiguous_cuda<float, nv_bfloat16>(src0_ddc, src1_ddc, ne, main_stream, dest_ptrs_d, graph_cpynode_index);
471-
} else {
472-
ggml_cpy_flt_cuda<float, nv_bfloat16> (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream, dest_ptrs_d, graph_cpynode_index);
473-
}
434+
ggml_cpy_flt_cuda<float, nv_bfloat16> (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream, dest_ptrs_d, graph_cpynode_index);
474435
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F16) {
475-
if (fast_cpy) {
476-
ggml_cpy_flt_contiguous_cuda<float, half>(src0_ddc, src1_ddc, ne, main_stream, dest_ptrs_d, graph_cpynode_index);
477-
} else {
478-
ggml_cpy_flt_cuda<float, half> (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream, dest_ptrs_d, graph_cpynode_index);
479-
}
436+
ggml_cpy_flt_cuda<float, half> (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream, dest_ptrs_d, graph_cpynode_index);
480437
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_Q8_0) {
481438
ggml_cpy_f32_q8_0_cuda(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream, dest_ptrs_d, graph_cpynode_index);
482439
} else if (src0->type == GGML_TYPE_Q8_0 && src1->type == GGML_TYPE_F32) {
@@ -548,7 +505,6 @@ void ggml_cuda_dup(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
548505
}
549506

550507
void* ggml_cuda_cpy_fn(const ggml_tensor * src0, ggml_tensor * src1) {
551-
bool fast_cpy = ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && ggml_are_same_shape(src0, src1);
552508
if (src0->type == src1->type && ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
553509
// Prioritize CUDA graph compatibility over direct memory copy optimization.
554510
// Using copy kernels here maintains graph indirection support, preventing performance regression from disabled CUDA graphs.
@@ -558,11 +514,11 @@ void* ggml_cuda_cpy_fn(const ggml_tensor * src0, ggml_tensor * src1) {
558514
return nullptr;
559515
}
560516
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32) {
561-
return fast_cpy ? (void *)cpy_flt_contiguous<float, float> : (void*) cpy_flt<cpy_1_flt<float, float>>;
517+
return (void*) cpy_flt<cpy_1_flt<float, float>>;
562518
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_BF16) {
563-
return fast_cpy ? (void *)cpy_flt_contiguous<float, nv_bfloat16> : (void*) cpy_flt<cpy_1_flt<float, nv_bfloat16>>;
519+
return (void*) cpy_flt<cpy_1_flt<float, nv_bfloat16>>;
564520
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F16) {
565-
return fast_cpy ? (void *)cpy_flt_contiguous<float, half> : (void*) cpy_flt<cpy_1_flt<float, half>>;
521+
return (void*) cpy_flt<cpy_1_flt<float, half>>;
566522
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_Q8_0) {
567523
return (void*) cpy_f32_q<cpy_blck_f32_q8_0, QK8_0>;
568524
} else if (src0->type == GGML_TYPE_Q8_0 && src1->type == GGML_TYPE_F32) {

‎ggml/src/ggml-cuda/mmvq.cu‎

Lines changed: 3 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -168,10 +168,9 @@ void ggml_cuda_op_mul_mat_vec_q_3D(
168168
GGML_UNUSED(src1_ddf_i);
169169
}
170170

171-
void ggml_cuda_op_mul_mat_vec_q_biased(
171+
void ggml_cuda_op_mul_mat_vec_q(
172172
ggml_backend_cuda_context & ctx,
173-
const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, const ggml_tensor * bias,
174-
const char * src0_dd_i, const float * src1_ddf_i,
173+
const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, const char * src0_dd_i, const float * src1_ddf_i,
175174
const char * src1_ddq_i, float * dst_dd_i, const int64_t row_low, const int64_t row_high, const int64_t src1_ncols,
176175
const int64_t src1_padded_row_size, cudaStream_t stream) {
177176

@@ -181,37 +180,14 @@ void ggml_cuda_op_mul_mat_vec_q_biased(
181180

182181
const int64_t ne0 = dst->ne[0];
183182

184-
if (bias) {
185-
if (bias->ne[0] != ne0) {
186-
printf("Oops: bias %s is %ld x %ld x %ld x %ld, dst %s is %ld x %ld x %ld x %ld\n",
187-
bias->name, bias->ne[0], bias->ne[1], bias->ne[2], bias->ne[3],
188-
dst->name, dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3]);
189-
}
190-
GGML_ASSERT(bias->ne[0] == ne0);
191-
GGML_ASSERT(bias->type == GGML_TYPE_F32);
192-
if (ggml_nrows(bias) != 1) {
193-
printf("Oops: bias %s is %ld x %ld x %ld x %ld\n", bias->name, bias->ne[0], bias->ne[1], bias->ne[2], bias->ne[3]);
194-
}
195-
GGML_ASSERT(ggml_nrows(bias) == 1);
196-
}
197-
198183
ggml_cuda_op_mul_mat_vec_q_impl(ctx, src0->type,
199184
ne00, ne0, 1, 0, 0, 0, 0, 0,
200-
src0_dd_i, nullptr, src1_ddq_i, dst_dd_i, nullptr, bias ? bias->data : nullptr, nullptr,
185+
src0_dd_i, nullptr, src1_ddq_i, dst_dd_i, nullptr, nullptr, nullptr,
201186
row_low, row_high, src1_ncols,
202187
src1_padded_row_size, GGML_UNARY_OP_COUNT, stream);
203188

204189
GGML_UNUSED(src1_ddf_i);
205190
}
206-
void ggml_cuda_op_mul_mat_vec_q(
207-
ggml_backend_cuda_context & ctx,
208-
const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst,
209-
const char * src0_dd_i, const float * src1_ddf_i,
210-
const char * src1_ddq_i, float * dst_dd_i, const int64_t row_low, const int64_t row_high, const int64_t src1_ncols,
211-
const int64_t src1_padded_row_size, cudaStream_t stream) {
212-
ggml_cuda_op_mul_mat_vec_q_biased(ctx, src0, src1, dst, nullptr, src0_dd_i, src1_ddf_i, src1_ddq_i, dst_dd_i, row_low, row_high, src1_ncols,
213-
src1_padded_row_size, stream);
214-
}
215191

216192
void ggml_cuda_op_mul_mat_vec_q_id(
217193
ggml_backend_cuda_context & ctx,

‎ggml/src/ggml-cuda/mmvq.cuh‎

Lines changed: 1 addition & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -9,20 +9,12 @@
99

1010
#define MMVQ_MAX_BATCH_SIZE 8 // Max. batch size for which to use MMVQ kernels.
1111

12-
void ggml_cuda_op_mul_mat_vec_q_biased(ggml_backend_cuda_context & ctx,
13-
const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, const ggml_tensor * bias,
14-
const char * src0_dd_i, const float * src1_ddf_i,
15-
const char * src1_ddq_i, float * dst_dd_i, const int64_t row_low, const int64_t row_high, const int64_t src1_ncols,
16-
const int64_t src1_padded_row_size, cudaStream_t stream);
17-
1812
void ggml_cuda_op_mul_mat_vec_q(ggml_backend_cuda_context & ctx,
19-
const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst,
20-
const char * src0_dd_i, const float * src1_ddf_i,
13+
const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, const char * src0_dd_i, const float * src1_ddf_i,
2114
const char * src1_ddq_i, float * dst_dd_i, const int64_t row_low, const int64_t row_high, const int64_t src1_ncols,
2215
const int64_t src1_padded_row_size, cudaStream_t stream);
2316

2417
bool ggml_cuda_mmvq_type_supported(ggml_type src0_type);
25-
2618
void ggml_cuda_op_mul_mat_vec_q_3D(ggml_backend_cuda_context & ctx,
2719
const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, const char * src0_dd_i, const float * src1_ddf_i,
2820
const char * src1_ddq_i, float * dst_dd_i, const int64_t row_low, const int64_t row_high, const int64_t src1_ncols,

‎src/llama-build-context.cpp‎

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1240,17 +1240,14 @@ std::tuple<ggml_tensor*, ggml_tensor*, ggml_tensor*> llm_build_context::llm_buil
12401240
if (bq) {
12411241
Qcur = ggml_add(ctx0, Qcur, bq);
12421242
cb(Qcur, "Qcur", il);
1243-
ggml_build_forward_expand(gf, Qcur);
12441243
}
12451244
if (bk) {
12461245
Kcur = ggml_add(ctx0, Kcur, bk);
12471246
cb(Kcur, "Kcur", il);
1248-
ggml_build_forward_expand(gf, Kcur);
12491247
}
12501248
if (bv) {
12511249
Vcur = ggml_add(ctx0, Vcur, bv);
12521250
cb(Vcur, "Vcur", il);
1253-
ggml_build_forward_expand(gf, Vcur);
12541251
}
12551252
return {Qcur, Kcur, Vcur};
12561253
}

0 commit comments

Comments
 (0)