NVIDIA CMP 50HX: Еще не все потеряно…

diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuhindex 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.cuindex 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 на dp4acmake -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 + эмуляция dp4acmake -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_DP4Afor n in a b c d e; do cmake --build build_$n -j --target llama-bench; done

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

llama.cpp

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

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

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

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

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

whisper.cpp

cd $WORKgit clone https://github.com/ggml-org/whisper.cppcd whisper.cpp# необязательно: git checkout 0c7be69git 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   # ожидается 1done

Если вы видите 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.cppcmake -B build_e "${CMP_FLAGS[@]}" -DLLAMA_CURL=OFFcmake --build build_e -j $JOBS --target llama-server llama-cli llama-bench

whisper.cpp

cd $WORK/whisper.cppcmake -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:" $fdoneldd $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 эту проблему решает, не убирайте его.