From f54cd74ed00a6095e6981d2ee6b4e801da6626a2 Mon Sep 17 00:00:00 2001 From: bssrdf Date: Sun, 16 Nov 2025 13:08:01 -0500 Subject: [PATCH] due to cp.async, only support filter size <= 32 --- ggml/src/ggml-cuda/conv2d-implicit.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/conv2d-implicit.cu b/ggml/src/ggml-cuda/conv2d-implicit.cu index 5958c3f29e..ce6c2b69d8 100644 --- a/ggml/src/ggml-cuda/conv2d-implicit.cu +++ b/ggml/src/ggml-cuda/conv2d-implicit.cu @@ -1309,7 +1309,7 @@ static void launch_conv2d_implicit_split_kernel(ggml_backend_cuda_context & ctx, static void conv2d_implicit_cuda_f16(ggml_backend_cuda_context & ctx, const float * X_D, const half * K_D, float * Y_D, int cc, param_t P, cudaStream_t st) { // if (GGML_CUDA_CC_IS_NVIDIA(cc) && turing_mma_available(cc) && P.c % 8 == 0 && (P.r > 1 || P.s > 1)) { - if (GGML_CUDA_CC_IS_NVIDIA(cc) && turing_mma_available(cc) && P.c % 8 == 0) { + if (GGML_CUDA_CC_IS_NVIDIA(cc) && turing_mma_available(cc) && P.c % 8 == 0 && (P.r <= 32 && P.s <= 32)) { int id = ggml_cuda_get_device();