Spaces:
Sleeping
Sleeping
Commit
·
6ccb5a5
1
Parent(s):
e9910b5
CUDA: fix matrix multiplication logic for tests (llama/6667)
Browse files- ggml-cuda.cu +1 -1
ggml-cuda.cu
CHANGED
|
@@ -1946,7 +1946,7 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor
|
|
| 1946 |
} else if (!split && !fp16_performance_good && src0->type == GGML_TYPE_F16 && !ggml_is_contiguous(src0) && !ggml_is_transposed(src1) && src1->ne[1] == 1) {
|
| 1947 |
// KQV single-batch
|
| 1948 |
ggml_cuda_mul_mat_vec_nc(ctx, src0, src1, dst);
|
| 1949 |
-
} else if (!split &&
|
| 1950 |
// KQ + KQV multi-batch
|
| 1951 |
ggml_cuda_mul_mat_batched_cublas(ctx, src0, src1, dst);
|
| 1952 |
} else if (use_dequantize_mul_mat_vec) {
|
|
|
|
| 1946 |
} else if (!split && !fp16_performance_good && src0->type == GGML_TYPE_F16 && !ggml_is_contiguous(src0) && !ggml_is_transposed(src1) && src1->ne[1] == 1) {
|
| 1947 |
// KQV single-batch
|
| 1948 |
ggml_cuda_mul_mat_vec_nc(ctx, src0, src1, dst);
|
| 1949 |
+
} else if (!split && src0->type == GGML_TYPE_F16 && (src1->type == GGML_TYPE_F16 || fp16_performance_good) && !ggml_is_transposed(src0) && !ggml_is_transposed(src1) && src1->ne[2]*src1->ne[3] > 1) {
|
| 1950 |
// KQ + KQV multi-batch
|
| 1951 |
ggml_cuda_mul_mat_batched_cublas(ctx, src0, src1, dst);
|
| 1952 |
} else if (use_dequantize_mul_mat_vec) {
|