NVIDIA CMP 50HX: шанс ещё остаётся…

diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh
index fea1cf9..046728a 100644
--- a/ggml/src/ggml-cuda/common.cuh
+++ b/ggml/src/ggml-cuda/common.cuh
@@ -738,7 +738,17 @@ static __device__ __forceinline__ int ggml_cuda_dp4a(const int a, const int b, i
 
 #else // defined(GGML_USE_HIP)
 
-#if __CUDA_ARCH__ >= GGML_CUDA_CC_DP4A || defined(GGML_USE_MUSA)
+#if defined(GGML_CUDA_CMP_EMU_DP4A) && !defined(GGML_USE_MUSA)
+    // NVIDIA CMP mining GPUs: the hardware dp4a is throttled (about 1/8 rate on CMP 40HX/50HX).
+    // Four plain integer multiply-adds written as inline PTX are faster and are not folded back
+    // into dp4a by the compiler (a plain C expression would be).
+    int r = c;
+    asm("mad.lo.s32 %0, %1, %2, %0;" : "+r"(r) : "r"((int)(int8_t)(a      )), "r"((int)(int8_t)(b      )));
+    asm("mad.lo.s32 %0, %1, %2, %0;" : "+r"(r) : "r"((int)(int8_t)(a >>  8)), "r"((int)(int8_t)(b >>  8)));
+    asm("mad.lo.s32 %0, %1, %2, %0;" : "+r"(r) : "r"((int)(int8_t)(a >> 16)), "r"((int)(int8_t)(b >> 16)));
+    asm("mad.lo.s32 %0, %1, %2, %0;" : "+r"(r) : "r"((int)(int8_t)(a >> 24)), "r"((int)(int8_t)(b >> 24)));
+    return r;
+#elif __CUDA_ARCH__ >= GGML_CUDA_CC_DP4A || defined(GGML_USE_MUSA)
     return __dp4a(a, b, c);
 #else // __CUDA_ARCH__ >= GGML_CUDA_CC_DP4A || defined(GGML_USE_MUSA)
     const int8_t * a8 = (const int8_t *) &a;
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index 9fe618f..c9ed96e 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -1403,6 +1403,36 @@ static __global__ void k_compute_batched_ptrs(
     ptrs_dst[0*ne23 + i12 + i13*ne12] = (      char *)         dst + i12*nbd2 + i13*nbd3;
 }
 
+#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
+// NVIDIA CMP mining GPUs (CMP 40HX/50HX/90HX): tensor cores and FMA are throttled, so cuBLAS must be
+// told not to use them (see https://habr.com/ru/articles/948396/: FP16 GEMM goes from ~3 to ~24
+// "units" with CUBLAS_COMPUTE_16F_PEDANTIC on a CMP 50HX).
+// GGML_CUDA_CMP_MODE=1 forces this on, =0 forces it off, unset/other = auto-detect by device name.
+static bool ggml_cuda_device_is_cmp(const int device) {
+    static int8_t cache[GGML_CUDA_MAX_DEVICES] = {0}; // 0 = unknown, 1 = no, 2 = yes
+    if (device < 0 || device >= GGML_CUDA_MAX_DEVICES) {
+        return false;
+    }
+    if (cache[device] == 0) {
+        const char * env = getenv("GGML_CUDA_CMP_MODE");
+        bool is_cmp;
+        if (env != nullptr && env[0] == '0') {
+            is_cmp = false;
+        } else if (env != nullptr && env[0] == '1') {
+            is_cmp = true;
+        } else {
+            cudaDeviceProp prop;
+            is_cmp = cudaGetDeviceProperties(&prop, device) == cudaSuccess && strstr(prop.name, "CMP") != nullptr;
+        }
+        cache[device] = is_cmp ? 2 : 1;
+        if (is_cmp) {
+            GGML_LOG_INFO("%s: device %d treated as NVIDIA CMP: cuBLAS FP16 GEMM without tensor cores (PEDANTIC)\n", __func__, device);
+        }
+    }
+    return cache[device] == 2;
+}
+#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
+
 // Type traits for mapping ggml types to CUDA/cuBLAS types
 template
 struct batched_mul_mat_traits;
@@ -1577,6 +1607,15 @@ static void ggml_cuda_mul_mat_cublas_impl(ggml_backend_cuda_context & ctx, const
     GGML_ASSERT(ne12 % ne02 == 0);
     GGML_ASSERT(ne13 % ne03 == 0);
 
+    auto gemm_algo = CUBLAS_GEMM_DEFAULT_TENSOR_OP;
+#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
+    if (compute_type == GGML_TYPE_F16 && cu_compute_type == CUBLAS_COMPUTE_16F && ggml_cuda_device_is_cmp(ctx.device)) {
+        // NVIDIA CMP: the tensor-core path is throttled; pure FP16 CUDA-core GEMM is ~7x faster.
+        cu_compute_type = CUBLAS_COMPUTE_16F_PEDANTIC;
+        gemm_algo       = CUBLAS_GEMM_DEFAULT;
+    }
+#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
+
     // broadcast factors
     const int64_t r2 = ne12/ne02;
     const int64_t r3 = ne13/ne03;
@@ -1599,7 +1638,7 @@ static void ggml_cuda_mul_mat_cublas_impl(ggml_backend_cuda_context & ctx, const
                            src1_ptr, cu_data_type_b, s11,
                     beta,   dst_ptr, cu_data_type,   ldc,
                     cu_compute_type,
-                    CUBLAS_GEMM_DEFAULT_TENSOR_OP));
+                    gemm_algo));
     } else if (r2 == 1 && r3 == 1 && is_src0_cont_2 && is_src1_cont_2) {
         // with a [0, 2, 1, 3] perm. and ne02==1 the matrix strides need to be determined from dim 3:
         const int64_t sma = ne02 == 1 ? s03 : s02;
@@ -1615,7 +1654,7 @@ static void ggml_cuda_mul_mat_cublas_impl(ggml_backend_cuda_context & ctx, const
                 beta,   dst_ptr, cu_data_type,   ldc, ne1*ldc, // strideC
                 ne12*ne13,
                 cu_compute_type,
-                CUBLAS_GEMM_DEFAULT_TENSOR_OP));
+                gemm_algo));
     } else {
         // use cublasGemmBatchedEx
         const int64_t ne23 = ne12*ne13;
@@ -1653,7 +1692,7 @@ static void ggml_cuda_mul_mat_cublas_impl(ggml_backend_cuda_context & ctx, const
                 beta,  (      void **) (ptrs_dst.get() + 0*ne23), cu_data_type,   ldc,
                 ne23,
                 cu_compute_type,
-                CUBLAS_GEMM_DEFAULT_TENSOR_OP));
+                gemm_algo));
     }
 
     // Convert output back to F32 if needed

# A: Базовая сборка (без патча)
cmake -B build_a -DGGML_CUDA=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_CUDA_COMPILER=/usr/local/cuda-12.8/bin/nvcc -DCMAKE_CUDA_ARCHITECTURES=75
# B: Код Pascal: без тензорных ядер, MMQ на dp4a
cmake -B build_b -DGGML_CUDA=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_CUDA_COMPILER=/usr/local/cuda-12.8/bin/nvcc -DCMAKE_CUDA_ARCHITECTURES="61-virtual" -DGGML_CUDA_FORCE_MMQ=ON
# C: B + запрет FMA во всех ядрах
cmake -B build_c -DGGML_CUDA=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_CUDA_COMPILER=/usr/local/cuda-12.8/bin/nvcc -DCMAKE_CUDA_ARCHITECTURES="61-virtual" -DGGML_CUDA_FORCE_MMQ=ON \
      -DCMAKE_CUDA_FLAGS="-fmad=false"
# D: C + эмуляция dp4a
cmake -B build_d -DGGML_CUDA=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_CUDA_COMPILER=/usr/local/cuda-12.8/bin/nvcc -DCMAKE_CUDA_ARCHITECTURES="61-virtual" -DGGML_CUDA_FORCE_MMQ=ON \
      -DCMAKE_CUDA_FLAGS="-fmad=false -DGGML_CUDA_CMP_EMU_DP4A"
# E: Гибрид: тензорные ядра для MMQ сохранены, FMA и dp4a убраны
cmake -B build_e -DGGML_CUDA=ON -DCMAKE_BUILD_TYPE=Release -DCMAKE_CUDA_COMPILER=/usr/local/cuda-12.8/bin/nvcc -DCMAKE_CUDA_ARCHITECTURES=75 -fmad=false -DGGML_CUDA_CMP_EMU_DP4A

for n in a b c d e; do cmake --build build_$n -j --target llama-bench; done