diff options
| -rw-r--r-- | ggml/src/ggml-cuda/mmvq.cu | 5 |
1 files changed, 4 insertions, 1 deletions
diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index d0477835..30a6a58b 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -72,10 +72,13 @@ static __device__ void mul_mat_vec_q( constexpr vec_dot_q_cuda_t vec_dot_q_cuda = get_vec_dot_q_cuda(type); + //int64_t rows_per_cuda_block = ggml_cuda_info().devices[id].cc < CC_RDNA2 ? + // ncols_y < 4 ? 1 : 2 : 1; + #if defined(GGML_USE_HIPBLAS) && defined(__HIP_PLATFORM_AMD__) && (defined(RDNA2) || defined(RDNA3)) constexpr int rows_per_cuda_block = 1; #else - constexpr int rows_per_cuda_block = ncols_y == 1 ? 1 : 2; + constexpr int rows_per_cuda_block = ncols_y < 4 ? 1 : 2; #endif // defined(GGML_USE_HIPBLAS) && defined(__HIP_PLATFORM_AMD__) && !defined(RDNA2) && !defined(RDNA3) const int tid = WARP_SIZE*threadIdx.y + threadIdx.x; |