Существует множество программных угроз для систем ИИ. Но я расскажу о GDDRHammer — атаке на аппаратную часть. Разберу эту атаку до винтика, то есть покажу, как она реализуется на физическом уровне, и сравню с атакой GeForge того же класса.

Атаки класса Rowhammer

Еще в 2014-м исследователи обнаружили, что, постоянно обращаясь к строкам оперативной памяти, могут изменить бит в соседних строках. Поскольку биты хранятся в виде электрического заряда, «взлом» их вызывает электрические помехи, способные превращать нули в единицы или наоборот. 

А через сбой в работе бита потенциальный злоумышленник может получить доступ к более высоким привилегиям в памяти, откуда скомпрометирует процессор. Этот класс атак получил название Rowhammer — по одной и той же строке (row) бьют много раз, словно молотком (hammer).

В 2025-м исследователи аналогично атаковали видеопамять. GPUHammer — первая успешная атака Rowhammer GDDR6 дискретных GPU NVIDIA. 

В отличие от классической Rowhammer на CPU, где атакующий «долбит» две строки памяти в надежде перевернуть бит в соседней строке, в GPUHammer три действия. 

Во-первых, исследователи провели обратный инжиниринг проприетарного драйвера NVIDIA и выявили функцию, раскрывающую физическое отображение виртуальной памяти на банки GDDR6. Это позволило точно выбрать в памяти строки, с помощью которых атаковали остальную память. Далее буду называть их строками-агрессорами. 

Во-вторых, авторы написали специализированное ядро CUDA, где тысячи параллельных потоков одновременно обращаются к одним и тем же строкам-агрессорам — параллелизм GPU создает лавину запросов, которую контроллер памяти со своим защитным механизмом Target Row Refresh не успевает обработать. 

В-третьих, атаку синхронизировали с циклами обновления памяти: «долбление» активируется сразу после того, как контроллер обновил строки, чтобы помехи накапливались максимально долго — до следующего обновления. 

В результате вызвали управляемые битовые искажения в таблицах страниц GPU, благодаря чему нарушили изоляцию памяти между процессами — первый шаг к полной компрометации системы. 

Тогда исследователи добились лишь восьми битовых искажений на всю видеокарту RTX A6000. Достаточно, чтобы продемонстрировать принцип и слегка подпортить работу нейросети, но недостаточно для серьезного взлома. Да, даже один битовый сбой может быть критичен, однако вероятность попасть в нужный бит за восемь переворотов стремится к нулю.

Атака GDDRHammer 

Авторы GPUHammer вернулись к задаче с новыми идеями и превзошли себя в 32 раза: в среднем 129 переворотов бита на каждый банк памяти RTX A6000. Но достижение не в самом количестве, а в том, что с ним атака из игры в рулетку становится инженерной задачей с предсказуемым результатом.

Что именно изменилось? Почему видеопамять, которая год назад казалась устойчивой к массированным атакам, демонстрирует уязвимость? Ответ в двух инновациях GDDRHammer: технике «подготовки памяти» (memory massaging) и использовании параллелизма GPU.

Memory massaging

Как я сказал, классический Rowhammer на CPU работает просто: выбираете две строки памяти (rows) в одном банке, «долбите» их тысячами обращений, и в какой-то момент электрические помехи переворачивают бит в соседней строке — жертве. 

Контроллеры памяти научились противостоять такой грубой силе. Их защищает технология Target Row Refresh (TRR). Механизм отслеживает частоту обращений к строкам. Когда счетчик обращений к определенной строке превышает порог, контроллер памяти принудительно обновляет соседние строки, «сбрасывая» накопленный электрический заряд и предотвращая переворот битов. Как если бы вы раскачивали качели, а кто-то каждые несколько секунд их останавливал.

GDDRHammer обходит TRR, уменьшая зону атаки. Суть техники memory massaging («подготовка памяти») — сдвинуть важные данные в области видеопамяти, наиболее уязвимые для электрических помех. Для этого GDDRHammer заполняет другие области своими данными, «подталкивая» критические структуры туда, где защита TRR работает хуже или где плотность ячеек эффективнее распространяет помехи. Атака использует штатные функции выделения памяти — например, cudaMalloc в экосистеме CUDA.

Еще исследователи выяснили, что данные в виртуальном адресном пространстве расположены не всегда так же, как физически на чипе. Memory massaging справляется и с этим: вслепую, без доступа к карте памяти, выявляет уязвимые области, измеряя время доступа и другие косвенные признаки.

Параллелизм как множитель силы

Второй секрет эффективности GDDRHammer — он использует параллельную природу GPU. 

В CPU атаку Rowhammer выполняют потоком или потоками, которые последовательно обращаются к строкам-агрессорам. Не быстро, даже если ядер несколько.

А графический процессор может одновременно выполнять тысячи потоков (warps в терминологии NVIDIA). И все они способны обращаться к памяти параллельно. Когда тысяча потоков одновременно «долбят» одни и те же строки памяти, эффект накапливается гораздо быстрее.

Исследователи GDDRHammer разработали специальное ядро CUDA, где тысячи потоков одновременно выполняют циклы чтения из двух тщательно выбранных строк. Код выглядит примерно так:

__global__ void hammer_kernel(volatile uint32_t *target) {
    // Каждый поток тысячи раз читает из двух соседних строк
    for (int i = 0; i < HAMMER_ITERATIONS; i++) {
        volatile uint32_t a = target[ROW_A_OFFSET];
        volatile uint32_t b = target[ROW_B_OFFSET];
        // Компилятор не должен оптимизировать эти операции
    }
}

Контроллер памяти GDDR6 не успевает обрабатывать лавину запросов и адекватно применять защитные механизмы. Частота обновления TRR рассчитана на определенную интенсивность обращений, а GPU-параллелизм легко превышает этот порог.

Результат подхода «memory massaging плюс параллельное „долбление“» — 129 битовых искажений на банк памяти. Достаточно, чтобы не просто случайно повредить данные, а целенаправленно модифицировать критические структуры управления памятью.

Основные цели атаки

Так почему 129 переворотов битов — это угроза? Потому что они позволяют атакующему добраться до таблиц страниц графического процессора. 

Таблицы страниц GPU — структуры данных, связывающие виртуальные адреса, видимые программами, с физическими адресами в видеопамяти. Когда ваше CUDA-приложение выделяет память через cudaMalloc, драйвер создает или обновляет записи в этих таблицах, сопоставляя виртуальный адрес, который вы получили, с физическим участком VRAM.

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

Допустим, старший бит физического адреса перевернулся с 0 на 1. Это изменит физический адрес, на который указывает страница, с сотен мегабайт на гигабайты. Теперь, когда атакующий обращается к своему виртуальному адресу, он на самом деле читает или пишет другую область видеопамяти — возможно, ту, где хранятся данные другой задачи или, что гораздо опаснее, структуры управления драйвера.

Исследователи GDDRHammer пошли дальше. Они показали, что за 129 попыток на банк можно не просто случайно повредить таблицу страниц, а контролируемо изменить конкретную запись. 

С помощью memory massaging они отправляют целевую запись таблицы страниц в физическую область, где вероятность битового искажения максимальна. Затем запускают параллельное «долбление» с точно рассчитанными параметрами (число итераций, выбор строк-агрессоров). После каждого цикла проверяют, произошел ли желаемый переворот. В случае неудачи — корректируют параметры и повторяют. За более чем сотню попыток вероятность успеха приближается к единице.

Когда таблица страниц повреждена, изоляция памяти рушится. Атакующий получает возможность читать и писать произвольные области видеопамяти, включая те, которые принадлежат другим процессам или самому драйверу. Уже серьезная брешь. 

Но GDDRHammer позволяет шагнуть и дальше — выбраться из видеопамяти в оперативную память центрального процессора. Здесь все становится еще интереснее.

Прорыв в CPU-память

Финальный и самый эффектный этап атаки GDDRHammer использует особенность того, как GPU взаимодействует с основной памятью. Современные GPU в режиме unified memory могут обращаться к оперативной памяти CPU через механизм DMA (Direct Memory Access). Для этого в таблицах страниц GPU существуют специальные записи, которые указывают не на видеопамять, а на области системной RAM.

GDDRHammer позволяет изменить запись в таблице страниц, которая указывала на легитимную область видеопамяти, таким образом, что она начинает указывать на память CPU. Например, меняется физический адрес в записи со значения в диапазоне VRAM на значение в диапазоне системной RAM.

После этого атакующее приложение, которое выполняется на GPU (помните, атака запускается через CUDA-код), получает возможность читать и писать произвольные области оперативной памяти центрального процессора. Злоумышленник может прочитать пароли, ключи шифрования, содержимое файлов и модифицировать критически важные структуры ядра операционной системы.

Один из исследователей, Эндрю Квонг, сформулировал это так: «Мы демонстрируем, как атакующий может вызвать битовый сбой на GPU, чтобы получить доступ к чтению и записи всей памяти CPU, что полностью компрометирует машину».

Технические детали реализации

Для полноты картины рассмотрю параметры, которые использовали исследователи. Атаку проводили на RTX A6000 — профессиональной видеокарте на архитектуре Ampere с 48 ГБ видеопамяти GDDR6. Ключевые условия: точный выбор строк-агрессоров (определенные пары строк в одном банке), число итераций «долбления» на поток (подбиралось экспериментально для обхода TRR), количество одновременно работающих потоков (тысячи), а также последовательность обращений к памяти внутри цикла.

Успешность атаки проверяли техникой «водяных знаков» — в потенциально уязвимые области памяти записали специальные паттерны и посмотрели, изменились ли они после «долбления». Таким образом не только детектировали перевороты, но и определили, какие биты затронуты.

Результаты воспроизводились стабильно. На RTX A6000 команда добивалась в среднем 129 флипов на банк, причем в некоторых банках число переворотов достигало нескольких сотен. Это означает, что, даже если несколько флипов пришлись на «неинтересные» области, атакующему хватит ресурсов поразить целевые структуры.

GDDRHammer vs GeForge

GDDRHammer — не единственная атака этого поколения. Параллельно с ней развивалась GeForge, которая атакует не таблицы страниц, а каталог таких таблиц (page directory) — структуру более высокого уровня, указывающую на таблицы. Повреждая каталог, можно скомпрометировать целый диапазон таблиц.

GeForge еще эффективнее на потребительских видеокартах: 1171 переворот бита на банк на RTX 3060. Однако на профессиональном ускорителе RTX A6000 эффективность падала до 202 флипов (все равно больше GDDRHammer). Почти шестикратное расхождение объясняется разницей в архитектуре памяти потребительских и профессиональных GPU.

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

Вкратце о защите от атак

Это тема для другой статьи, здесь лишь обозначу тезисно. Сейчас нужно включать ECC (автокоррекцию ошибок) и IOMMU (диспетчер доступа к физическим участкам памяти). А в перспективе — менять архитектуру, делая будущие поколения GPU устойчивее к атакам класса Rowhammer.