Comments 15
Работает это так: пробуем последовательно инвертировать все подозрительные биты в калькуляторе, затем получаем новое шестнадцатеричное значение для байтов
Откуда берутся подозрительные биты в калькуляторе? Можете привести какой-нибудь пример?
4001400520008043
100000000000001010000000000010100100000000000001000000001000011
ISCADD R2.CC, R0, c[0x0][0x140], 0x2;
4001400500009C43
100000000000001010000000000010100000000000000001001110001000011
Индексы считаем для строк с двоичной записью числа, справа налево.
Для этих двух инструкций различия только в битах с индексами:
1) 10-12 — в первом случае значение 0, т.к. используется предикат P0, во втором — 7, т.к. предикат не используется, это зарезервированное значение;
2) 29 — здесь нужно попробовать декодировать шестнадцатеричные числа 148 и 140, тогда получим 101001000 и 101000000.
Видим, что начало в бите 26, конец определяем так: инвертируем все ведущие нули по очереди, переводя число каждый раз в шестнадцатеричную систему счисления, заменяя изменившиеся цифры/буквы в hex-редакторе для бинарного файла axpy.cubin, не забывая про little endian.
Затем сверяем вывод дизассемблера в каждом отдельном случае. После каждой попытки возвращаем биты в исходное состояние, меняя таким образом только 1 бит за раз. Так наткнёмся на лимит на 41 бите.
Результат декомпиляции
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;
}
А этот результат декомпиляции уже из гидры получен?
Но вообще написать реализацию на Pcode — задача даже более простая, чем писать грамматику для декодера байтов. Быстро получалось исправлять реализацию для некоторых сложных инструкций из x86 (и не только), благодаря очень удобному промежуточному языку, единому мидлэнду (оптимизатор), 2 бэкэндам (в основном C; как альтернативный вариант — Java/C#, больше похоже на последний, т.к. время от времени появляется goto, но не labeled break).
А что за middlend и backend имеется ввиду? Это же не гидровский?
И причем тут Java и C#?)
Как и в компиляторах, в декомпиляторах тоже есть middleend и backend. Middleend — оптимизации и внутренняя обработка. Backend — вывод на каком-то одном языке. По умолчанию, т.е. для машинных языков, используется C. Для некоторых, вроде JVM (байткод) и Dalvik, используется вывод на Java, но из-за goto напоминает C#, т.к. в Java вместо goto надо использовать break/continue. Но вообще это принято называть Java-подобным кодом, т.к. это на самом деле может быть и HLSL. Таким образом в Ghidra идёт разделение на языки машинные и управляемые.
Eclipse можно и не использовать, я использую Intellij Idea. На LLVM действительно чем-то похоже. К примеру, ни один человек не догадался объединить все языки при помощи IR/IL, в итоге имели кучу компиляторов в 2000-х годах, притом все — такое себе… С Ghidra то же самое. Всякие иды только и делали, что переводили байты в абстрактные заголовки без какой-либо реализации инструкций. В Ghidra реализация есть на Pcode, так что декомпилятор может оттранслировать байты не только в дизассемблерные заголовки (вроде XOR EAX,EAX), но и в микрокод. На счёт остальных инструментов — только если radare2/cutter с его ESIL. Микрокод можно также использовать, если не знаете ассемблер. К примеру я не вижу особого смысла запоминать длинющий список из мнемоник, тем более что был один забавный случай. Для FPU инструкций из x86 (вроде 64 битная версия) никогда в дизассемблере не указывается, что меняется регистр ST0, т.е. он может хранить возвращаемое значение. Поэтому, увидев почти пустую функцию в декомпиляторе я, не задумываясь, сразу начал смотреть текст Pcode, там и увидел ST0, а затем на википедии нашёл, что это входит в одно из соглашений вызовов для FPU.
Ну да, я тоже чуть не начал писать свой декомпилятор для байткода DirectX, чтобы автоматически декомпилировать шейдеры из игр Unity/Unreal. Но потом я быстро отказался от этой затеи, теперь вот планирую написать спецификацию-фронтенд для Ghidra. Но Ghidra — это не компилятор с ассемблером, а наоборот — декомпилятор с дизассемблером. Там всё проще ввиду отсутствия информации об ООП (в которое CPU/GPU всё равно не умеют), за исключением байткодов, где свой немножко подход. Плюс бэкэнд и миддлэнд везде один, т.е. не нужно реализовывать "схлопывание" нескольких инструкций IR в одно при помощи селектора и прочее. Если можно что-то поджать, например, какое-то выражение, то оно будет поджато одинаково независимо от архитектуры, если только конечный пользователь что-то не поменял в преференсах программы, но не плагина.
Всякие иды только и делали, что переводили байты в абстрактные заголовки без какой-либо реализации инструкций
Похоже что уже нет, вот интересная презентация https://binary.ninja/presentations/Modern%20Binary%20Analysis%20with%20ILs.pdf
судя по всему промежуточное предстaвление подвезли в IDA https://i.blackhat.com/us-18/Thu-August-9/us-18-Guilfanov-Decompiler-Internals-Microcode-wp.pdf
На счёт остальных инструментов — только если radare2/cutter с его ESIL.
В остальных тоже много где есть:
Но у Гидры Pcode, самый читаемый по сравнению с ESIL и Microcode(IDA). Однако не такой простой как у binary ninja BNIL, судя по их презентации))
Ну как такого BNIL нет, есть только целая серия языков, вот ссылку нашёл для ознакомления:
https://docs.binary.ninja/dev/bnil-llil.html
К примеру, в LILL есть ненужное разделение на single, double. В Pcode есть только длина операндов (регистр либо локальная варнода, аналог скрытых микрорегистров), этим все подобные проблемы решаются сразу. Плюс ненужные операции (т.н. синтаксический сахар), к примеру, сдвиги с переносом. Для этого в Ghidra используется развёрнутое примитивное выражение. Так, в x86 инструкция "ROR rm8,CL" имеет такую примитивную (простейшую, основанную на операциях, а не функциях) реализацию:
local cnt = CL & 0x7; # Остаток от деления на 8
rm8 = (rm8 >> cnt) # Младшая часть, основная
|| (rm8 << (8 — cnt)); # Старшая часть, переносная
В ESIL тоже есть подобное, плюс ещё свои операции. К примеру, equal-операции (++, --, +=, -=, *=, /=, %=). Так что в Pcode минимализм и примитивизм по сути. Остальное через бэкэнд (к примеру, inplace assignment для equal-операций, но не всегда срабатывает, иногда так и оставляет uVar1 = uVar1 + uVar2 вместо uVar1 += uVar2). Но длина Pcode реализаций зачастую получается больше за счёт урезанного алфавита-перечня допустимых операций. Единственное, в Pcode присутствуют функции вроде FLOAT_SQRT, FLOAT_ABS, но они считаются за blackbox. А вот syscall нет, вместо него CALLIND с возможностью приставить референс на функцию, являющейся представителем системного вызова, наверное, обработчиком (handler).
Пишем спецификацию под Nvidia Kepler (бинарники CUDA, версия языка sm_30) для Ghidra