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

Предыстория

Есть домашний сервер обычной “подстольной” конфигурации. Proxmox с несколькими виртуалками. В какой-то момент прикупил на аукционе NVIDIA CMP 50HX за копейки, ради поиграться с простенькими модельками. Модели запускаю через llama.cpp, дополнительно использую Whisper для распознавания речи. Быстро смирился со скоростью работы и жил с этим всем спокойно несколько месяцев, пока не наткнулся на статью NVIDIA CMP – микроскопы для забивания гвоздей? Копаем глубже….

Очень рекомендую прочесть статью, ощутимо просветляет мозг в части понимания проблемы (и варианта решения), о которых пойдет речь ниже.

Обращение к “высшему разуму”

Отправил ссылку на статью “в чат” с просьбой предложить решение. В итоге через 20 минут у меня был готовый стенд с патчем для исходников ggml. Поскольку ggml в llama.cpp и whisper.cpp общий, один и тот же патч подходит для обоих проектов.

Если процесс тестирования вам не интересен, можно сразу перейти к результатам и готовому решению.

llamacpp-cmp.patch

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

Важное замечание: Патч собран для актуальных (на дату публикации статьи) веток репозиториев. На llama.cpp (ветка a0e6dc7) патч накладывается чисто (git apply). На whisper.cpp (ветка 0c7be69) строки сдвинуты на 45, поэтому git apply отказывает, но patch -p1 применяет все пять фрагментов (fuzz 1–2). После наложения патча все три вызова GEMM в ggml_cuda_mul_mat_cublas_impl используют новую переменную gemm_algo.
Для будущих веток, очевидно патч нужно будет проверять (или переделывать)

Что делает патч

1. cuBLAS FP16 без тензорных ядер (ggml-cuda.cu)

В ggml_cuda_mul_mat_cublas_impl добавлена функция ggml_cuda_device_is_cmp(). Она определяет карту CMP по имени устройства и кэширует результат. На CMP вызовы cublasGemmEx / cublasGemmStridedBatchedEx / cublasGemmBatchedEx получают другие параметры:

CUBLAS_COMPUTE_16F → CUBLAS_COMPUTE_16F_PEDANTIC;
CUBLAS_GEMM_DEFAULT_TENSOR_OP → CUBLAS_GEMM_DEFAULT.

Поведением управляет переменная окружения GGML_CUDA_CMP_MODE: 1 принудительно включает режим, 0 выключает, без переменной работает автоопределение. При срабатывании в лог пишется строка:

ggml_cuda_device_is_cmp: device 0 treated as NVIDIA CMP: cuBLAS FP16 GEMM without tensor cores (PEDANTIC)

Переключается способ матричного умножения с весами в F16 (это прежде всего Whisper).

2. Эмуляция dp4a (common.cuh)

Если собрать с -DGGML_CUDA_CMP_EMU_DP4A, функция ggml_cuda_dp4a вместо __dp4a выполняет четыре инструкции mad.lo.s32, записанные Inline-PTX:

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;

Inline-PTX здесь нужен потому, что обычное C-выражение компилятор свернул бы обратно в dp4a. На dp4a построены ядра квантованного умножения матрицы на вектор (MMVQ) и MMQ без тензорных ядер, а значит, прежде всего генерация токенов.

Подготовка тестового стенда и сборка бенчмарков

Если процесс тестирования вам не интересен, можно сразу перейти к результатам и готовому решению.

0. Мой стенд


GPU	NVIDIA CMP 50HX, 10 ГБ, compute capability 7.5
Драйвер	580.126.09
CUDA Toolkit	12.8 (сборка), драйвер поддерживает CUDA 13.0
CPU / RAM	12 потоков, 32 ГБ
ОС	Proxmox VE (ядро 6.8.12-pve)
llama.cpp	`a0e6dc7` (master на 15.10.2026)
whisper.cpp	`0c7be69` (master на 15.10.2026)

1. Проверка оборудования и драйвера

nvidia-smi --query-gpu=name,compute_cap,driver_version,memory.total --format=csv

В колонке name должно быть слово CMP (например, NVIDIA CMP 50HX), иначе автоопределение в патче не сработает. Драйвер нужен с поддержкой CUDA 12.x: в выводе nvidia-smi должно быть CUDA Version: 12.x или выше.

Для других CMP подставьте свою архитектуру вместо 75:

Карта	Чип	`CMAKE_CUDA_ARCHITECTURES`
CMP 30HX, 40HX, 50HX	Turing	`75`

Проверялась только CMP 50HX

2. Пакеты и CUDA Toolkit

Debian / Ubuntu:

sudo apt update
sudo apt install -y build-essential cmake git ccache curl python3
cmake --version    # нужен 3.18+, проверено на 3.25

Убедитесь, что в выбранном тулките есть nvcc и cuBLAS:

ls $CUDA_ROOT/bin/nvcc
ls $CUDA_ROOT/include/cublas_v2.h $CUDA_ROOT/lib64/libcublas.so*
$CUDA_ROOT/bin/nvcc --version

Если cublas_v2.h нет, этот тулкит установлен неполностью; возьмите другой (на тестовом сервере в 12.8). Поставить cuBLAS можно так:

sudo apt install -y cuda-toolkit-12-8     # или libcublas-dev-12-8 при наличии репозитория NVIDIA

Посмотрите, куда указывает симлинк /usr/local/cuda:

readlink -f /usr/local/cuda

Если он ведёт не на $CUDA_ROOT (например, на CUDA 13), это не страшно, но в cmake обязательно передавайте -DCUDAToolkit_ROOT, иначе возможна ошибка линковки cudaGetDeviceProperties_v2.

3. Cборка бенчмарков для тестирования

Для проверки всех вариантов, собрал 5 вариантов с различной комбинацией опций и прогнал тестами. Каждый вариант собирался в отдельном каталоге с общими параметрами:

-DGGML_CUDA=ON \
-DCMAKE_BUILD_TYPE=Release \
-DCMAKE_CUDA_COMPILER=/usr/local/cuda-12.8/bin/nvcc

В таблице дополнительные параметры для тестирования различных вариантов сборки

Вариант	Архитектура	Доп. параметры	Сборка
A	75	—	Базовая сборка (без патча)
B	61-virtual	`GGML_CUDA_FORCE_MMQ=ON`	Код Pascal: без тензорных ядер, MMQ на dp4a
C	61-virtual	B + `-fmad=false`	B + запрет FMA во всех ядрах
D	61-virtual	C + `-DGGML_CUDA_CMP_EMU_DP4A`	C + эмуляция dp4a
E	75	`-fmad=false -DGGML_CUDA_CMP_EMU_DP4A`	Гибрид: тензорные ядра для MMQ сохранены, FMA и dp4a убраны

Команды для сборки

# 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

Результаты: llama.cpp

Методика замеров

LLM: llama-bench -m <модель> -p 512 -n 128 -ngl 99, по 5 повторов (значение по умолчанию), pp512 — скорость обработки промпта из 512 токенов (prefill), tg128 — скорость генерации 128 токенов. Обе метрики в токенах в секунду, больше — лучше.
Флаги запуска: -fa 1 -ctk q8_0 -ctv q8_0 (flash attention и KV-кэш в q8_0).
Whisper: whisper-cli с моделью large-v3-turbo на samples/jfk.wav (11 секунд аудио), время кодирования (encode) и общее время.
Очистка: Перед каждым замером GPU был пуст (0 МиБ занято), каждый прогон шёл примерно 5 минут.

1. Qwen3.5-9B Q4_K_M (5,28 ГиБ)

Сборка	pp512, ток/с	tg128, ток/с
A — базовая	307.8 ± 0.6	29.9 ± 0.5
B — Pascal + MMQ	84.4 ± 0.1	28.8 ± 0.6
C — B + без FMA	111.7 ± 0.1	44.7 ± 0.3
D — C + эмуляция dp4a	105.9 ± 0.1	54.6 ± 0.2
E — sm_75 + без FMA + эмуляция dp4a	412.0 ± 1.6	58.8 ± 0.2

2. Qwen3.8-9B-Distill Q4_K_M (5,37 ГиБ)

Сборка	pp512, ток/с	tg128, ток/с
A — базовая	309.5 ± 0.8	30.0 ± 0.7
B — Pascal + MMQ	84.7 ± 0.1	29.4 ± 0.6
C — B + без FMA	111.7 ± 0.1	45.5 ± 0.2
D — C + эмуляция dp4a	108.0 ± 0.1	55.1 ± 0.1
E — sm_75 + без FMA + эмуляция dp4a	412.6 ± 1.4	57.8 ± 0.1

3. Qwen3.8-9B-Distill Q5_K_M (6,18 ГиБ)

Сборка	pp512, ток/с	tg128, ток/с
A — базовая	311.4 ± 0.3	29.9 ± 0.2
B — Pascal + MMQ	84.7 ± 0.1	28.4 ± 0.7
C — B + без FMA	111.6 ± 0.1	44.2 ± 0.3
D — C + эмуляция dp4a	117.1 ± 0.1	51.2 ± 0.0
E — sm_75 + без FMA + эмуляция dp4a	410.7 ± 1.4	53.2 ± 0.2

Что показали замеры

1. Тензорные ядра для INT8 на CMP 50HX не урезаны, по крайней мере в MMQ. Варианты B–D, где тензорных ядер нет вообще, обрабатывают промпт в 3–4 раза медленнее базового. Значит обходить их не нужно.

2. Генерацию тормозят FMA и dp4a. Генерация токенов упирается в ядра MMVQ, а они построены на dp4a и FMA. Если сравнивать в рамках «Pascal-ветки», переход от B к C (запрет FMA) даёт +55% к генерации, а от C к D (эмуляция dp4a) ещё +16–23%.

3. Отсюда вариант E: лучшее от обоих. Оставляем sm_75, чтобы MMQ шёл через тензорные ядра, и убираем FMA и dp4a. Результат на Q4_K_M:

	pp512	tg128
A (было)	306.8	28.4
E (стало)	412.6	57.8
Ускорение	×1.34	×2.04

Отдельно стоит пояснить, почему E обгоняет A даже по prefill. Флаг -fmad=false действует на все ядра: на нормализации, softmax, RoPE, деквантизацию и flash attention. На CMP все они упираются в урезанный FMA.

4. Чем крупнее квант, тем ниже генерация, а prefill не меняется. Генерация упирается в объём весов: 58 ток/с на Q4_K_M против 53 на Q5_K_M. Prefill во всех случаях около 410 ток/с.

Результаты: whisper.cpp

Модель large-v3-turbo хранит веса в F16, поэтому здесь основную роль играет первая часть патча (cuBLAS PEDANTIC). Её эффект измерялся отдельно на сборке A, с переключением GGML_CUDA_CMP_MODE.

Сборка	encode, мс	total, мс
A, `GGML_CUDA_CMP_MODE=0`	873.7	2071.8
A, автоопределение (PEDANTIC включён)	265.7	1038.7
B	1097.1	2454.8
C	251.5	999.4
D	251.5	999.4
E	267.3	1020.6

Выводы:

PEDANTIC FP16 GEMM ускоряет энкодер в 3,3 раза (874 → 266 мс), а общее время сокращается вдвое. Прирост в 7 раз из статьи, измеренный на чистом GEMM, не получается, но в реальной задаче эффект всё равно большой.

Итоги

Задача	Было	Стало	Ускорение
LLM, генерация (tg128)	28.4 ток/с	57.8 ток/с	×2.0
LLM, обработка промпта (pp512)	306.8 ток/с	412.6 ток/с	×1.34
Whisper, энкодер (jfk.wav)	874 мс	~266 мс	×3.3

Что понятно из результатов тестирования CMP 50HX:

FMA действительно урезан, и это главный тормоз. Сборка с -fmad=false — самый дешёвый и самый эффективный шаг: код не меняется, нужен только флаг компилятора.
dp4a урезан. Эмуляция через замену четырех функций в mad.lo.s32 даёт ещё около +20% к генерации.
cuBLAS FP16 урезан. Режим PEDANTIC ускоряет F16-модели вроде whisper в 3 с лишним раза.
INT8 MMQ на тензорных ядрах работает нормально. Сборка под Pascal (61-virtual) снижает prefill втрое, поэтому делать её не нужно.

Рекомендуемая сборка для CMP 50HX (и, вероятно, для других CMP на Turing):

cmake -B build -DGGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=75 \
      -DCMAKE_CUDA_FLAGS="-fmad=false -DGGML_CUDA_CMP_EMU_DP4A"

плюс патч для cuBLAS PEDANTIC.

Инструкция для сборки

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 эту проблему решает, не убирайте его.

Выводы

Гипотеза, предложенная в исходной статье, вполне подтверждается.

Можно ли еще что-то выжать из этих карточек (думаю, что нет). Есть ли задачи для которых сейчас обоснованно использовать такие карты – решать вам. В моем случае выигрыш в скорости работы Whisper’a оказался весьма ощутимым.

Всем добра!

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

Предыстория

Обращение к “высшему разуму”

Что делает патч

1. cuBLAS FP16 без тензорных ядер (ggml-cuda.cu)

2. Эмуляция dp4a (common.cuh)

Подготовка тестового стенда и сборка бенчмарков

0. Мой стенд

1. Проверка оборудования и драйвера

2. Пакеты и CUDA Toolkit

3. Cборка бенчмарков для тестирования

Результаты: llama.cpp

Методика замеров

1. Qwen3.5-9B Q4_K_M (5,28 ГиБ)

2. Qwen3.8-9B-Distill Q4_K_M (5,37 ГиБ)

3. Qwen3.8-9B-Distill Q5_K_M (6,18 ГиБ)

Что показали замеры

Результаты: whisper.cpp

Выводы:

Итоги

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

llama.cpp

whisper.cpp

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

2. Сборка

llama.cpp

whisper.cpp

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

Выводы

Разработчик Telegram: 99,9% клиентского кода мессенджера пишет AI, а тех кто до сих пор пишет код вручную надо увольнять

Sony отобрала у PS3 Linux, а через год потеряла ключи от всей консоли

НКЦКИ при ФСБ потребовал срочно обновить PHP. Тем, кто не успевает, предложено отключить сайты

Почему корней DNS тринадцать, хотя в 512 байт влезало пятнадцать

Публикации