Для Ghidra уже написано довольно много спецификаций для обычных процессорных языков, но ничего для графических.
Это и понятно, ведь у него есть своя специфика: предикаты, константы, через которые передаются параметры, среди прочего, и прочее, унаследованное от шейдеров.
Кроме того, формат, используемый для хранения кода, часто является проприетарным, и вам придется изменить его самостоятельно.
В этой статье мы на двух примерах разберемся, что к чему.
Первая программа — это простейшая axpy (аналог hello world для GPGPU).
Второй помогает понять реализацию условий и прыгает на GPU, потому что там всё по-другому.
Все языки Nvidia используют кодировку с прямым порядком байтов, поэтому сразу копируем байты из шестнадцатеричного редактора в какой-нибудь блокнот (например, Notepad++) в обратном порядке по 8 штук (длина инструкций здесь постоянная).
Затем с помощью калькулятора программирования (подойдет тот, что от Microsoft) преобразуем его в двоичный код. Далее ищем совпадения, создаем маску инструкций, затем операндов.
Для декодирования и поиска маски использовались hex-редактор и дизассемблер cuobjdump; иногда требуется ассемблер, как в AMDGPU (поскольку дизассемблера там нет, но это тема отдельной статьи).
Работает это так: пытаемся последовательно инвертировать все подозрительные биты в калькуляторе, затем получаем новое шестнадцатеричное значение байт, подставляем их в бинарник, скомпилированный через nvcc или ассемблер, если он есть, что не всегда является правильным.
случай.
Затем проверяем с помощью cuobjdump. Выкладываю исходники в формате (в основном на C, без плюсов и ООП для более тесной связи с машинным GPU-кодом), потом сразу дизассем + байты, потому что так удобнее, их не надо менять местами.
Скопируйте в axpy.cu и скомпилируйте через cmd: nvcc axpy.cu --cubin --gpu-architecture sm_30 Результирующий ELF-файл с именем axpy.cubin дизассемблируется там: cuobjdump axpy.cubin -sass Пример 1:
Свалка__global__ void axpy(float param_1, float* param_2, float* param_3) { unsigned int uVar1 = threadIdx.x; param_2[uVar1] = param_1 * param_3[uVar1]; }
/*0000*/
/* 0x22c04282c2804307 */
/*0008*/ MOV R1, c[0x0][0x44];
/* 0x2800400110005de4 */
/*0010*/ S2R R0, SR_TID.X;
/* 0x2c00000084001c04 */
/*0018*/ MOV32I R5, 0x4;
/* 0x1800000010015de2 */
/*0020*/ ISCADD R2.CC, R0, c[0x0][0x150], 0x2;
/* 0x4001400540009c43 */
/*0030*/ LD.E R2, [R2];
/* 0x8400000000209c85 */
/*0038*/ ISCADD R4.CC, R0, c[0x0][0x148], 0x2;
/* 0x4001400520011c43 */
/*0040*/
/* 0x20000002e04283f7 */
/*0048*/ IMAD.U32.U32.HI.X R5, R0, R5, c[0x0][0x14c];
/* 0x208a800530015c43 */
/*0050*/ FMUL R0, R2, c[0x0][0x140];
/* 0x5800400500201c00 */
/*0058*/ ST.E [R4], R0;
/* 0x9400000000401c85 */
/*0060*/ EXIT;
/* 0x8000000000001de7 */
/*0068*/ BRA 0x68;
/* 0x4003ffffe0001de7 */
/*0070*/ NOP;
/* 0x4000000000001de4 */
/*0078*/ NOP;
/* 0x4000000000001de4 */Результат декомпиляции
void axpy(float param_1,float *param_2,float *param_3) {
uint uVar1;
uVar1 = *&threadIdx.x;
param_2[uVar1] = param_3[uVar1] * param_1;
return;
}Пример 2:
__global__ void predicates(float* param_1, float* param_2) {
unsigned int uVar1 = threadIdx.x + blockIdx.x * blockDim.x;
if ((uVar1 > 5) & (uVar1 < 10)) param_1[uVar1] = uVar1;
else param_2[uVar1] = uVar1;
}Свалка
/*0000*/
/* 0x2272028042823307 */
/*0008*/ MOV R1, c[0x0][0x44];
/* 0x2800400110005de4 */
/*0010*/ S2R R0, SR_TID.X;
/* 0x2c00000084001c04 */
/*0018*/ S2R R3, SR_CTAID.X;
/* 0x2c0000009400dc04 */
/*0020*/ IMAD R0, R3, c[0x0][0x28], R0;
/* 0x20004000a0301ca3 */
/*0028*/ MOV32I R3, 0x4;
/* 0x180000001000dde2 */
/*0030*/ IADD32I R2, R0, -0x6;
/* 0x0bffffffe8009c02 */
/*0038*/ I2F.F32.U32 R4, R0;
/* 0x1800000001211c04 */
/*0040*/
/* 0x22c042e04282c2c7 */
/*0048*/ ISETP.GE.U32.AND P0, PT, R2, 0x4, PT;
/* 0x1b0ec0001021dc03 */
/*0050*/ @P0 ISCADD R2.CC, R0, c[0x0][0x148], 0x2;
/* 0x4001400520008043 */
/*0058*/ @P0 IMAD.U32.U32.HI.X R3, R0, R3, c[0x0][0x14c];
/* 0x208680053000c043 */
/*0060*/ @P0 ST.E [R2], R4;
/* 0x9400000000210085 */
/*0068*/ @P0 EXIT;
/* 0x80000000000001e7 */
/*0070*/ ISCADD R2.CC, R0, c[0x0][0x140], 0x2;
/* 0x4001400500009c43 */
/*0078*/ MOV32I R3, 0x4;
/* 0x180000001000dde2 */
/*0080*/
/* 0x2000000002e04287 */
/*0088*/ IMAD.U32.U32.HI.X R3, R0, R3, c[0x0][0x144];
/* 0x208680051000dc43 */
/*0090*/ ST.E [R2], R4;
/* 0x9400000000211c85 */
/*0098*/ EXIT;
/* 0x8000000000001de7 */
/*00a0*/ BRA 0xa0;
/* 0x4003ffffe0001de7 */
/*00a8*/ NOP;
/* 0x4000000000001de4 */
/*00b0*/ NOP;
/* 0x4000000000001de4 */
/*00b8*/ NOP;
/* 0x4000000000001de4 */Результат декомпиляции
void predicates(float *param_1,float *param_2) {
uint uVar1;
uVar1 = *&blockIdx.x * (int)_DAT_constants_00000028 + *&threadIdx.x;
if (uVar1 - 6 < 4) {
param_1[uVar1] = (float)uVar1;
return;
}
param_2[uVar1] = (float)uVar1;
return;
}Нетрудно догадаться, что тесты изначально были рассчитаны на машинный код, чтобы компилятору нечего было оптимизировать.
Во всем остальном вам придется вручную отменить оптимизацию.
В сложных примерах это может быть вообще невозможно, поэтому в таких случаях придется довериться декомпилятору и фронтенду.
В общем правило такое: для тестирования фронтенда берём любой простой (с минимумом возможных оптимизаций) первый подходящий (воспроизводящий ошибки) пример.
В остальном декомпилированный код уже будет иметь оптимизации (или будет только как-то исправлен через рефакторинг).
Но пока главная задача — это хотя бы просто правильный код, который делает то же самое, что и машинный код. Это «Программное моделирование».
«Моделирование программного обеспечения» само по себе не включает в себя рефакторинг, перевод C в C++, восстановление классов и уж точно не такие вещи, как выявление шаблонов.
Теперь ищем шаблоны мнемоники, операндов и модификаторов.
Для этого мы сравниваем биты (в двоичном представлении) между подозрительными инструкциями (или строками, если их так удобнее называть).
Вы также можете использовать то, что другие пользователи публикуют в своих вопросах на stackoverflow, например «помогите мне понять двоичный/sass/машинный код», использовать учебные пособия (в том числе на китайском языке) и другие ресурсы.
Так, основной номер операции хранится в битах 58-63, но есть и дополнительные биты 0-4 (они различают инструкции «I2F», «ISETP», «MOV32I»), где-то вместо них 0-2 (чтобы пренебрегайте 3-4 битами в пустых инструкциях, в спецификации они отмечены как «UNK»).
По регистрам и постоянным числам можно поэкспериментировать дизассемблером, чтобы найти все биты, влияющие на вывод дампа типа того, что выложен под спойлером.
Все поля, которые мне удалось найти, есть в спецификации на Github, файл CUDA.slasspec, раздел токенов.
Потом нужно придумать адреса для регистров, они опять же есть на Github. Это необходимо, т.к.
на микроуровне Sleigh трактует регистры как глобальные переменные в пространстве с типом «register_space», но поскольку их пространство не помечено как «выводимое» (и уж точно не может быть), то в декомпиляторе они становятся либо локальные переменные (чаще всего с приставкой «Var», но иногда вроде бы присутствовал приставка «local»), либо параметры (префикс «param_»).
SP никогда не пригодился, он был нужен в основном формально для обеспечения функциональности декомпилятора нужен ПК (что-то вроде IP от x86) для эмуляции.
Затем есть регистры-предикаты, что-то вроде флагов, но более «общего назначения», чем для заранее заданной цели, например, переполнение, (не) равное нулю и т. д. Затем блокирующий регистр для имитации кучи инструкций ISCADD.CC и IMAD.HI, т.к.
первая из них в моей реализации выполняет вычисления за себя, а за вторую, чтобы не переносить часть суммы в старшие 4 байта, т.к.
это испортит декомпиляцию.
Но тогда вам нужно заблокировать следующий регистр до завершения операции IMAD.HI. Что-то подобное, т.е.
несоответствие официальной документации ожидаемому выводу декомпилятора уже было в модуле SPU для той же Гидры.
Кроме того, существуют специальные регистры, которые в настоящее время реализуются через cpool. В будущем я планирую заменить их символами, определенными по умолчанию для некоторого «выводимого» пространства.
Это одни и те же threadIdx, blockIdx. Затем привязываем переменные к полям dest, par0, par1, par2, res. Затем идут подтаблицы, а после них — то, ради чего все и затевалось — основные (корневые) таблицы с основными инструкциями.
Здесь необходимо строго соблюдать формат «мнемоника-операнды», но есть послабление для модификаторов, которые, тем не менее, должны быть прикреплены к мнемосхеме или к секции операндов.
Никакие другие форматы неприемлемы, даже ту же ассемблировку Hexagon DSP придется адаптировать под этот синтаксис, что, впрочем, не очень сложно.
Завершающим этапом будет написание реализации инструкций на языке микропрограммирования Pcode. Единственное, что хотелось бы отметить из первого примера, это инструкции ISCADD.CC и IMAD.HI, где первая из них берет указатель на регистры и разыменовывает их как указатели на 8 байт вместо 4. Это сделано намеренно в чтобы лучше приспособиться к декомпилятору и его поведению, несмотря на то, что в документации Nvidia написано о передаче части суммы.
Для второго примера лучше в настройках декомпилятора поставить галочку напротив «Упростить предикацию».
Дело в том, что предикаты — это одно и то же условие для разных инструкций, по сути не что иное, как всем известный «SIMD», или его следующий аналог.
Те.
если бит предиката установлен, то инструкции выполняются, причем подряд. Еще нужно завести привычку сразу писать реализацию для каждой инструкции, а не только прототип (мнемоника-операнд), ведь есть еще декомпилятор, эмулятор и прочие анализаторы.
Но в целом написание реализации на Pcode — ещё более простая задача, чем написание грамматики для байтового декодера.
Быстро удалось подправить реализацию некоторых сложных инструкций от x86 (и не только), благодаря очень удобному промежуточному языку, единому мидленду (оптимизатору), 2-м бэкендам (в основном C; как альтернатива - Java/C#, подробнее как и последнее, т. е.
потому что время от времени появляется goto, но не помечается как перерыв).
В будущих статьях также могут быть интерфейсы для управляемых языков, таких как DXBC, SPIR-V, они будут использовать бэкенд Java/C#.
Но пока в планах только машинные коды, потому что.
байткоды требуют особого подхода.
Проект Гидра Запросы: Pкод Сани Теги: #nvidia #C++ #обратное проектирование #GPGPU #Ghidra #декомпиляция #pcode #sleigh #трансляция кода
-
Распространение Ultimate Edition 2.6
19 Oct, 24 -
Где Спрятана История Космоса?
19 Oct, 24 -
Нетскейп?
19 Oct, 24
