В этой статье на двух примерах разберёмся, что к чему.
Первая программа простейший axpy (аналог hello world для GPGPU). Вторая помогает понять реализацию условий и прыжков на GPU, т.к. там всё по-другому.
Во всех Nvidia языках используется кодировка little endian, так что сразу копируем байты из hex-редактора в какой-нибудь блокнот (например, 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;}
Нетрудно догадаться, что тесты изначально заточены под машинный код, чтобы компилятору нечего было оптимизировать. Для всего остального пришлось бы вручную отменять оптимизации. В сложных же примерах это вовсе может быть невозможно, так что для таких случаев придётся довериться декомпилятору и фронтенду.
Вообще правило такое для тестирования фронтенда берём любой простой (с минимумом возможных оптимизаций) первый подходящих (воспроизводящий ошибки) пример. Для остального декомпилированный код уже будет с оптимизациями (либо только как-то через рефакторинг поправлять). Но пока что основная задача хотя бы просто верный код, делающий то же самое, что и машинный. Это и есть Software modelling. Сам Software modelling не предполагает рефакторинг, перевод C в C++, восстановление классов, и уж тем более таких вещей как идентификация шаблонов.
Теперь ищем паттерны для мнемоники, операндов и модификаторов.
Для этого сравниваем биты (в двоичном представлении) между подозрительными инструкциями (или строками, если их так удобнее называть). Можно также воспользоваться тем, что выкладывают другие пользователи в своих вопросах на stackoverflow по типу помогите понять двоичный/sass/машинный код, задействовать туториалы (в т.ч. на китайском языке) и прочие ресурсы. Так, основной номер операции хранится в битах 58-63, но есть и дополнительные биты 0-4 (они различают инструкции I2F, ISETP, MOV32I), где-то вместо них 0-2 (для пренебрежения 3-4 битами в пустых инструкциях, в спецификации они отмечены как UNK).
Для регистров и чисел-констант можно экспериментировать с дизассемблером дабы найти все биты, влияющие на вывод дампа наподобие выложенного под спойлером. Все поля, которые мне удалось найти, есть в спецификации на Github, файл CUDA.slaspec, раздел token.
Затем нужно придумать адреса для регистров, опять же они есть на Github. Это нужно, т.к. на микроуровне Sleigh рассматривает регистры как глобальные переменные в пространстве с типом register_space, но т.к. их пространство не отмечено как inferable (и наверняка оно не может быть), то они в декомпиляторе становятся либо локальными переменными (чаще всего с интерфиксом Var, но иногда вроде был и префикс local), либо параметрами (префикс param_). SP так и не пригодился, нужен в основном формально для обеспечения работоспособности декомпилятора. PC (что-то вроде IP из x86) нужен для эмуляции.
Затем идут предикатовые регистры, что-то вроде флагов, но уже более general purpose, чем для заранее продиктованной цели, вроде переполнения, (не)равенства нулю и т.п.
Затем блокировочный регистр для моделирования связки инструкций ISCADD .CC и IMAD.HI, т.к. первая из них в моей реализации выполняет подсчет за себя и за вторую, чтобы избежать переноса части суммы в старшие 4 байта, т.к. это испортит декомпиляцию. Но тогда надо заблокировать следующий регистр до завершения операции IMAD.HI. Что-то подобное, т.е. разночтение официальной документации и ожидаемого вывода декомпилятора, уже было в модуле SPU для той же Ghidra.
Потом идут специальные регистры, которые пока что реализованы через cpool. В будущем я планирую их заменить на символы, определённые по умолчанию для какого нибудь inferable пространства. Это те самые threadIdx, blockIdx.
Затем привязываем переменные к полям dest, par0, par1, par2, res. Затем идут подтаблицы, а после них то, ради чего всё и затевалось основные (корневые) таблицы с главными инструкциями.
Здесь нужно строго следовать формату мнемоника-операнды, однако даётся послабление для модификаторов, которые, тем не менее, должны быть прикреплены к мнемонике либо к секции с операндами. Никакие другие форматы недопустимы, даже тот же Hexagon DSP asm придётся адаптировать к этому синтаксису, что впрочем не очень сложно.
Финальным этапом будет написание реализации для инструкций на языке микропрограммирования Pcode. Единственное, что хотелось бы отметить из первого примера, это инструкции ISCADD .CC и IMAD.HI, где первая из них берёт указатель на регистры и разыменовывает их как указатели на 8 байтов вместо 4. Это сделано намеренно для того, чтобы лучше приспособиться к декомпилятору и его поведению, несмотря на то, что написано в документации Nvidia про перенос части суммы.
Для второго примера лучше поставить галочку в настройках декомпилятора напротив надписи Simplify predication. Смысл в том, что предикаты это одно и то же условие для разных инструкций, по сути ни что иное как всем известное SIMD, либо его очередной аналог. Т.е. если предикатовый бит установлен, то инструкции выполняются, притом подряд.
Нужно также завести привычку сразу писать реализацию для каждой инструкции, а не только прототип (мнемоника-операнды), ведь есть ещё декомпилятор, эмулятор и другие анализаторы.
Но вообще написать реализацию на Pcode задача даже более простая, чем писать грамматику для декодера байтов. Быстро получалось исправлять реализацию для некоторых сложных инструкций из x86 (и не только), благодаря очень удобному промежуточному языку, единому мидлэнду (оптимизатор), 2 бэкэндам (в основном C; как альтернативный вариант Java/C#, больше похоже на последний, т.к. время от времени появляется goto, но не labeled break).
В следующих статьях, возможно, будут также фронтенды для managed языков, таких как DXBC, SPIR-V, они будут использовать бэкэнд Java/C#. Но пока что в планах только машинные коды, т.к. байткоды требуют особого подхода.
Проект
Ghidra
Справки:
Pcode
Sleigh