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<ggml_type T>
 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

1. Исходники и патч

llama.cpp

mkdir -p $WORK && cd $WORK
git clone https://github.com/ggml-org/llama.cpp
cd llama.cpp
# необязательно: зафиксировать проверенную версию
git checkout a0e6dc7
git apply --check -v $PATCH && git apply $PATCH
git diff --stat

Должно получиться:

 ggml/src/ggml-cuda/common.cuh   | 12 ++++++++++-
 ggml/src/ggml-cuda/ggml-cuda.cu | 45 ++++++++++++++++++++++++++++++++++++++---

Если git apply отказывает (у вас другая версия), попробуйте по очереди:

git apply --3way $PATCH                 # только для полного клона, не --depth 1
patch -p1 --dry-run < $PATCH && patch -p1 < $PATCH

whisper.cpp

cd $WORK
git clone https://github.com/ggml-org/whisper.cpp
cd whisper.cpp
# необязательно: git checkout 0c7be69
git apply --check $PATCH && git apply $PATCH 
  || (patch -p1 --dry-run < $PATCH && patch -p1 < $PATCH)

В whisper.cpp патч обычно ложится через patch со сдвигом строк и fuzz. Это нормально, но после этого проверьте результат.

Проверка наложения патча (для обоих проектов)

for r in llama.cpp whisper.cpp; do
  echo "== $r"
  f=$WORK/$r/ggml/src/ggml-cuda/ggml-cuda.cu
  grep -c "gemm_algo));" $f                         # ожидается 3
  grep -c "CUBLAS_GEMM_DEFAULT_TENSOR_OP));" $f     # ожидается 0
  grep -c "GGML_CUDA_CMP_EMU_DP4A" $WORK/$r/ggml/src/ggml-cuda/common.cuh   # ожидается 1
done

Если вы видите CUBLAS_GEMM_DEFAULT_TENSOR_OP)); где-то ещё, значит, в вашей версии есть вызовы cuBLAS, которые патч не покрывает. Замените в них CUBLAS_GEMM_DEFAULT_TENSOR_OP на gemm_algo вручную по образцу и добавьте тот же блок if (... ggml_cuda_device_is_cmp(...)).

2. Сборка

Общие параметры:

CMP_FLAGS=(
  -DCUDAToolkit_ROOT=$CUDA_ROOT
  -DCMAKE_CUDA_COMPILER=$CUDA_ROOT/bin/nvcc
  -DGGML_CUDA=ON
  -DCMAKE_BUILD_TYPE=Release
  -DBUILD_SHARED_LIBS=OFF
  -DCMAKE_CUDA_ARCHITECTURES=75
  "-DCMAKE_CUDA_FLAGS=-fmad=false -DGGML_CUDA_CMP_EMU_DP4A"
)
export PATH=$CUDA_ROOT/bin:$PATH

llama.cpp

cd $WORK/llama.cpp
cmake -B build_e "${CMP_FLAGS[@]}" -DLLAMA_CURL=OFF
cmake --build build_e -j $JOBS --target llama-server llama-cli llama-bench

whisper.cpp

cd $WORK/whisper.cpp
cmake -B build_e "${CMP_FLAGS[@]}"
cmake --build build_e -j $JOBS --target whisper-server whisper-cli whisper-bench

Сборка занимает 5–10 минут на проект, в зависимости от CPU.

Проверка, что собралось правильно

for f in $WORK/llama.cpp/build_e/CMakeCache.txt $WORK/whisper.cpp/build_e/CMakeCache.txt; do
  grep -E "^CMAKE_CUDA_FLAGS:|^CMAKE_CUDA_ARCHITECTURES|^CUDA_cudart_LIBRARY:" $f
done
ldd $WORK/llama.cpp/build_e/bin/llama-server | grep -E "cublas|cudart"

Во всех путях должен стоять $CUDA_ROOT (например, cuda-12.8), а не /usr/local/cuda → 13.x.

Сборка от root. sudo сбрасывает PATH, поэтому export PATH выше не действует. Флаг -DCUDAToolkit_ROOT в CMP_FLAGS эту проблему решает, не убирайте его.