На протяжении более чем пятнадцати лет процессоры архитектуры x86 поставлялись с выделенной аппаратной инструкцией для умножения без переноса (carryless multiplication). Это кажется незначительной деталью, но на самом деле этот примитив лежит в основе аутентифицированного шифрования, кодов, исправляющих ошибки, и современных систем нулевого разглашения (Zero-Knowledge Proofs). Долгое время графические процессоры NVIDIA оставались в стороне от этой эволюции, заставляя разработчиков использовать обходные пути, такие как битовые срезы (bitslicing), которые были эффективны, но не оптимальны.
С выходом CUDA 13.3 эта ситуация кардинально меняется. NVIDIA внедрила инструкцию clmad (carryless multiply-accumulate) для всех GPU на базе архитектуры Ampere и новее (SM 80+). Это не просто очередное обновление библиотеки; это фундаментальное изменение в том, как GPU обрабатывает арифметику в конечных полях. В этой статье мы подробно разберем, как работает эта инструкция, почему она так важна для кибербезопасности и AI, и как разработчики могут интегрировать её в свои проекты уже сегодня.
01Почему GPU не мог умножать «без переноса» раньше?
Чтобы понять масштаб прорыва, нужно заглянуть в математику. Криптографические алгоритмы часто оперируют не обычными целыми числами, а элементами расширенных бинарных полей (binary extension fields). Простейшее конечное поле GF(2) состоит из одного бита, где сложение — это XOR, а умножение — AND. Однако для реальной криптографии этого мало. Мы используем поля GF(2^n), где n битов представляют коэффициенты полинома над GF(2).
Сложение в таких полях остается простым побитовым XOR. Но умножение требует «длинного умножения» битов каждого входа, с последующим приведением по модулю неприводимого полинома. Например, в поле GF(2^2) с неприводимым полиномом x^2 + x + 1, результат умножения двух элементов должен быть корректно приведен.
До появления clmad разработчикам на GPU приходилось эмулировать эту операцию через серию инструкций AND, XOR и сдвигов. Это создавало значительные накладные расходы. Каждый шаг эмуляции требовал извлечения отдельных битов, маскирования и сдвига, что забивало исполнительные блоки GPU и увеличивало задержку. На CPU эта проблема была решена давно благодаря инструкции PCLMULQDQ, но на GPU она оставалась «узким горлышком» для криптографических workload'ов.
02Что такое clmad и как она работает?
Инструкция clmad — это аппаратно ускоренная команда умножения без переноса с накоплением. Она доступна в виде инструкции PTX (Parallel Thread Execution) в CUDA 13.3. По своей сути она аналогична PCLMULQDQ на x86, но адаптирована для архитектуры NVIDIA GPU.
Ключевая особенность clmad заключается в том, что она выполняет умножение двух входных значений, давая расширенный результат, и одновременно добавляет аккумулятор. Это критически важно для полиномиальной арифметики, где промежуточные результаты часто нужно суммировать «на лету» без потери точности.
Инструкция имеет две вариации:
clmad.lo.u64— вычисляет нижние биты результата.clmad.hi.u64— вычисляет верхние биты результата.
Это позволяет разработчикам работать с большими числами, разбивая их на части, что идеально подходит для операций в полях GF(2^128) и GF(2^256).
clmad работает напрямую с регистрами общего назначения (gpr) в PTX. Это означает, что вы можете использовать её в inline-ассемблере внутри CUDA-ядер, что дает максимальный контроль над производительностью, но требует понимания архитектуры PTX.Пример кода: Умножение больших чисел
Ниже приведен пример того, как можно реализовать умножение больших чисел, используя clmad. Этот код демонстрирует, как разбить задачу на две операции и собрать результат.
__device__ inline uint128_t clmad_mul_128(uint64_t a, uint64_t b, uint128_t acc) {
uint64_t acc_lo = (uint64_t)acc;
uint64_t acc_hi = (uint64_t)(acc >> 64);
uint64_t lo, hi;
// Inline PTX: calculate lower and higher bits of result
asm("clmad.lo.u64 %0, %1, %2, %3;" : "=l"(lo) : "l"(a), "l"(b), "l"(acc_lo));
asm("clmad.hi.u64 %0, %1, %2, %3;" : "=l"(hi) : "l"(a), "l"(b), "l"(acc_hi));
return ((uint128_t)hi << 64) | lo;
}Этот подход позволяет избежать дорогостоящих сдвигов и маскирования, которые требовались в предыдущих реализациях на базе ALU.
03Ускорение GHASH: Сердце AES-GCM
Одним из самых важных применений clmad является ускорение алгоритма GHASH, который является частью стандарта аутентифицированного шифрования AES-GCM. AES-GCM используется повсеместно: в TLS для защиты веб-трафика, в VPN, в системах хранения данных и в большинстве корпоративных решений для шифрования «в покое» (at rest) и «в движении» (in transit).
GHASH вычисляет хеш аутентифицирования для всех входных данных (шифротекста, дополнительных аутентифицированных данных и блока длины). Он разбивает входные данные на блоки, XOR-ит каждый блок с текущим аккумулятором и умножает результат на ключ хеширования H (полученный путем шифрования нулевого блока с помощью AES) в поле GF(2^128).
Именно это умножение в поле GF(2^128) было самым медленным компонентом. С clmad оно выполняется значительно быстрее, обеспечивая пропускную способность, сопоставимую с пропускной способностью памяти DRAM, и достигая ускорения почти в 19 раз по сравнению с предыдущими методами на базе битовых срезов на NVIDIA B200.
Результаты бенчмарков
Команда NVIDIA провела масштабное тестирование производительности GHASH на двух ключевых платформах: NVIDIA GeForce RTX 5090 (для энтузиастов и рабочих станций) и NVIDIA B200 (для дата-центров и AI-класт
Источник: NVIDIA Developer ↗