diff --git a/ggml/src/ggml-cuda.cu b/ggml/src/ggml-cuda.cu
index 10959122..cfc9caa9 100644
--- a/ggml/src/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda.cu
@@ -1604,33 +1604,6 @@ static void ggml_cuda_op_mul_mat_cublas(
         }
         const half * src1_ptr = src1->type == GGML_TYPE_F16 ? (const half *) src1_ddf_i : src1_as_f16.get();
 
-        // On Volta, avoid storing f32 graph outputs in a temporary f16 buffer;
-        // finite matmul results outside fp16 range would become +/-inf there.
-        const bool sm70_f32_output =
-            compute_capability <= CC_VOLTA &&
-            dst->type == GGML_TYPE_F32;
-        if (sm70_f32_output) {
-            const float alpha_f32 = 1.0f;
-            const float beta_f32  = 0.0f;
-
-            static std::atomic<int> sm70_f32_output_logs{0};
-            if (sm70_f32_output_logs.fetch_add(1) < 8) {
-                GGML_CUDA_LOG_WARN(
-                        "%s: using f32 cublas output for %s on cc=%d to avoid fp16 output saturation\n",
-                        __func__, dst->name, compute_capability);
-            }
-
-            CUBLAS_CHECK(cublasSetStream(ctx.cublas_handle(id), stream));
-            CUBLAS_CHECK(
-                cublasGemmEx(ctx.cublas_handle(id), CUBLAS_OP_T, CUBLAS_OP_N,
-                        row_diff, src1_ncols, ne10,
-                        &alpha_f32, src0_ptr,  CUDA_R_16F, ne00,
-                                    src1_ptr,  CUDA_R_16F, ne10,
-                        &beta_f32,  dst_dd_i,  CUDA_R_32F, ldc,
-                        CUBLAS_COMPUTE_32F,
-                        CUBLAS_GEMM_DEFAULT_TENSOR_OP));
-            return;
-        }
         ggml_cuda_pool_alloc<half> dst_f16(ctx.pool(id), row_diff*src1_ncols);
 
         const half alpha_f16 = 1.0f;
diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu
index fde49334..5991547e 100644
--- a/ggml/src/ggml-cuda/mmq.cu
+++ b/ggml/src/ggml-cuda/mmq.cu
@@ -247,7 +247,9 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11) {
 #endif //GGML_CUDA_FORCE_MMQ
 
     if (cc < CC_OFFSET_AMD) {
-        return cc < CC_VOLTA || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;
+        // On Volta, large-batch quantized matmuls otherwise fall back through
+        // fp16 cuBLAS temporaries. Keep using MMQ for pre-Turing NVIDIA.
+        return cc < CC_TURING || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;
     }
 
     return cc < CC_RDNA3 || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;
diff --git a/ggml/src/ggml-cuda/mmq_id.cu b/ggml/src/ggml-cuda/mmq_id.cu
index 1f3b2f6f..13bab3dc 100644
--- a/ggml/src/ggml-cuda/mmq_id.cu
+++ b/ggml/src/ggml-cuda/mmq_id.cu
@@ -566,7 +566,9 @@ bool ggml_cuda_can_use_mmq_id(enum ggml_type type, int cc, int64_t ne11) {
 #endif //GGML_CUDA_FORCE_MMQ
 
     if (GGML_CUDA_CC_IS_NVIDIA(cc)) {
-        return !fp16_mma_hardware_available(cc) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;
+        // Match the plain MMQ policy: use MMQ for pre-Turing NVIDIA, including
+        // Volta, so indexed/expert matmuls avoid the fp16 cuBLAS fallback.
+        return cc < CC_TURING || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;
     }
 
    if (amd_mfma_available(cc)) {