Русский
Русский
English
Статистика
Реклама

Декомпиляция

Меняем промежуточное представление кода на лету в Ghidra

30.04.2021 14:04:59 | Автор: admin

Когда мы разрабатывали модуль ghidra nodejs для инструмента Ghidra, мы поняли, что не всегда получается корректно реализовать опкод V8 (движка JavaScript, используемого Node.js) на языке описания ассемблерных инструкций SLEIGH. В таких средах исполнения, как V8, JVM и прочие, один опкод может выполнять достаточно сложные действия. Для решения этой проблемы в Ghidra предусмотрен механизм динамической инъекции конструкций P-code языка промежуточного представления Ghidra. Используя этот механизм, нам удалось превратить вывод декомпилятора из такого:

В такой:

Рассмотрим пример с опкодом CallRuntime. Он вызывает одну функцию из списка т.н. Runtime-функций V8 по индексу (kRuntimeId). Также данная инструкция имеет переменное число аргументов (range номер начального регистра-аргумента, rangedst число аргументов). Описание инструкции на языке SLEIGH, который Ghidra использует для определения ассемблерных инструкций, выглядит так:

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

  1. Поиск нужного названия функции в массиве Runtime-функций по индексу kRuntimeId.

  2. Поскольку аргументы передаются через регистры, необходимо сохранить их предыдущее состояние.

  3. Передача в функцию переменного количества аргументов.

  4. Вызов функции и сохранение результата вызова в аккумулятор.

  5. Восстановление предыдущего состояния регистров.

Если вы знаете, как сделать такое на SLEIGH, пожалуйста, напишите комментарий. А мы решили, что все это (а особенно работу с переменным количеством аргументов-регистров) не очень удобно (если возможно) реализовывать на языке описания процессорных инструкций, и применили механизм динамических инъекций p-code, который как раз для таких случаев реализовали разработчики Ghidra. Что это за механизм?

Можно создать в файле описания ассемблерных инструкций (slaspec) специальную пользовательскую операцию, например CallRuntimeCallOther. Далее, изменив конфигурацию вашего модуля (подробнее об этом ниже), вы можете сделать так, чтобы при нахождении в коде данной инструкции Ghidra передавала бы обработку в Java динамически, и уже на языке Java написать обработчик, который будет динамически формировать p-code для инструкции, пользуясь всей гибкостью Java.

Рассмотрим подробно, как это сделать.

Создание служебной операции SLEIGH

Опишем опкод CallRuntime следующим образом. Подробнее об описании процессорных инструкций на языке SLEIGH все можете узнать из статьи Создаем процессорный модуль под Ghidra на примере байткода v8.

Определим служебную операцию:

define pcodeop CallRuntimeCallOther;

И опишем саму инструкцию:

:CallRuntime [kRuntimeId], range^rangedst is op = 0x53; kRuntimeId; range;       rangedst {CallRuntimeCallOther(2, 0);}

Таким образом, любой опкод, начинающийся с байта 0x53, будет расшифрован как CallRuntime При попытке его декомпиляции будет вызываться обработчик операции CallRuntimeCallOtherс аргументами 2 и 0. Эти аргументы описывают тип инструкции (CallRuntime) и позволят нам написать один обработчик для нескольких похожих инструкций (CallWithSpread, CallUndefinedReceiverи т.п.).

Подготовительная работа

Добавим класс, через который будет проходить инъекция кода: V8_PcodeInjectLibrary. Этот класс мы унаследуем от ghidra.program.model.lang.PcodeInjectLibrary который реализует большую часть необходимых для инъекции p-code методов.

Начнем написание класса V8_PcodeInjectLibraryс такого шаблона:

package v8_bytecode;import public class V8_PcodeInjectLibrary extends PcodeInjectLibrary {public V8_PcodeInjectLibrary(SleighLanguage l) {}}

V8_PcodeInjectLibraryбудет использоваться не пользовательским кодом, а движком Ghidra, поэтому нам необходимо задать значение параметра pcodeInjectLibraryClassв файле pspec, чтобы движок Ghidra знал, какой класс задействовать для инъекции p-code.

<?xml version="1.0" encoding="UTF-8"?><processor_spec>  <programcounter register="pc"/>  <properties>  <property key="pcodeInjectLibraryClass" value="v8_bytecode.V8_PcodeInjectLibrary"/>  </properties></processor_spec>

Также нам понадобится добавить нашу инструкцию CallRuntimeCallOtherв файл cspec. Ghidra будет вызывать V8_PcodeInjectLibraryтолько для инструкций, определенных таким образом в cspec-файле.

<callotherfixup targetop="CallRuntimeCallOther"><pcode dynamic="true"><input name=outsize"/> </pcode></callotherfixup>

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

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

public class V8_PcodeInjectLibrary extends PcodeInjectLibrary {private Set<String> implementedOps;private SleighLanguage language;public V8_PcodeInjectLibrary(SleighLanguage l) {super(l);language = l;String translateSpec = language.buildTranslatorTag(language.getAddressFactory(),getUniqueBase(), language.getSymbolTable());PcodeParser parser = null;try {parser = new PcodeParser(translateSpec);}catch (JDOMException e1) {e1.printStackTrace();}implementedOps = new HashSet<>();implementedOps.add("CallRuntimeCallOther");}}

Благодаря внесенным нами изменениям Ghidra будет вызывать метод getPayloadнашего класса V8_PcodeInjectLibraryкаждый раз при попытке декомпиляции инструкции CallRuntimeCallOther Создадим данный метод, который при наличии инструкции в списке реализованных операций будет создавать объект класса V8_InjectCallVariadic(этот класс мы реализуем чуть позже) и возвращать его.

@Override/*** This method is called by DecompileCallback.getPcodeInject.*/public InjectPayload getPayload(int type, String name, Program program, String context) {if (type == InjectPayload.CALLMECHANISM_TYPE) {return null;}if (!implementedOps.contains(name)) {return super.getPayload(type, name, program, context);}V8_InjectPayload payload = null; switch (name) {case ("CallRuntimeCallOther"):payload = new V8_InjectCallVariadic("", language, 0);break;default:return super.getPayload(type, name, program, context);}return payload;}

Генерация p-code

Основная работа по динамическому созданию p-code будет происходить в классе V8_InjectCallVariadic. Давайте его создадим и опишем типы операций.

package v8_bytecode;import public class V8_InjectCallVariadic extends V8_InjectPayload {public V8_InjectCallVariadic(String sourceName, SleighLanguage language, long uniqBase) {super(sourceName, language, uniqBase);}// Типы операций. В данном примере мы рассматриваем RUNTIMETYPEint INTRINSICTYPE = 1;int RUNTIMETYPE = 2;int PROPERTYTYPE = 3;@Overridepublic PcodeOp[] getPcode(Program program, InjectContext context) {}@Overridepublic String getName() {return "InjectCallVariadic";}}

Как нетрудно догадаться, нам необходимо разработать нашу реализацию метода getPcode Для начала создадим объект pCode класса V8_PcodeOpEmitter Этот класс будет помогать нам создавать инструкции pCode (позже мы ознакомимся с ним подробнее).

V8_PcodeOpEmitter pCode = new V8_PcodeOpEmitter(language, context.baseAddr, uniqueBase); 

Далее из аргумента context (контекст инъекции кода) мы можем получить адрес инструкции, который нам пригодится в дальнейшем.

Address opAddr = context.baseAddr;

С помощью данного адреса мы получим объект текущей инструкции:

Instruction instruction = program.getListing().getInstructionAt(opAddr);

Также с помощью аргумента contextмы получим значения аргументов, которые ранее описывали на языке SLEIGH.

Integer funcType = (int) context.inputlist.get(0).getOffset();Integer receiver = (int) context.inputlist.get(1).getOffset();

Реализуем обработку инструкции и генерации Pcode.

// проверка типа инструкцииif (funcType != PROPERTYTYPE) {// получаем kRuntimeId  индекс вызываемой функцииInteger index = (int) instruction.getScalar(0).getValue();// сгенерируем Pcode для вызова инструкции cpool с помощью объекта pCode класса V8_PcodeOpEmitter. Подробнее остановимся на нем ниже.pCode.emitAssignVarnodeFromPcodeOpCall("call_target", 4, "cpool", "0", "0x" + opAddr.toString(), index.toString(), funcType.toString());}// получаем аргумент диапазон регистровObject[] tOpObjects = instruction.getOpObjects(2);// get caller args count to save only necessary onesObject[] opObjects;Register recvOp = null;if (receiver == 1) {}else {opObjects = new Object[tOpObjects.length];System.arraycopy(tOpObjects, 0, opObjects, 0, tOpObjects.length);}// получаем количество аргументов вызываемой функцииtry {callerParamsCount = program.getListing().getFunctionContaining(opAddr).getParameterCount();}catch(Exception e) {callerParamsCount = 0;}// сохраняем старые значения регистров вида aN на стеке. Это необходимо для того, чтобы Ghidra лучше распознавала количество аргументов вызываемой функцииInteger callerArgIndex = 0;for (; callerArgIndex < callerParamsCount; callerArgIndex++) {pCode.emitPushCat1Value("a" + callerArgIndex);}// сохраняем аргументы вызываемой функции в регистры вида aNInteger argIndex = opObjects.length;for (Object o: opObjects) {argIndex--;Register currentOp = (Register)o;pCode.emitAssignVarnodeFromVarnode("a" + argIndex, currentOp.toString(), 4);}// вызов функцииpCode.emitVarnodeCall("call_target", 4);// восстанавливаем старые значения регистров со стекаwhile (callerArgIndex > 0) {callerArgIndex--;pCode.emitPopCat1Value("a" + callerArgIndex);}// возвращаем массив P-Code операцийreturn pCode.getPcodeOps();

Теперь рассмотрим логику работы класса V8_PcodeOpEmitter (https://github.com/PositiveTechnologies/ghidra_nodejs/blob/main/src/main/java/v8_bytecode/V8_PcodeOpEmitter.java), который во многом основан на аналогичном классе модуля для JVM. Данный класс генерирует p-code операции с помощью ряда методов. Рассмотрим их в порядке обращения к ним в нашем коде.

emitAssignVarnodeFromPcodeOpCall(String varnodeName, int size, String pcodeop, String... args)

Для понимания работы данного метода сначала рассмотрим понятие Varnodeодин из основных элементов p-code, по сути представляющий собой любую переменную, задействованную в p-code. Регистры, локальные переменные всё это Varnode.

Вернемся к методу. Данный метод генерирует p-code для вызова функции pcodeopс аргументами argsи сохраняет результат работы функции в varnodeName То есть в итоге получается такая конструкция:

varnodeName = pcodeop(args[0], args[1], );

emitPushCat1Value(String valueName) и emitPopCat1Value (String valueName)

Генерирует p-code для аналогов ассемблерных операций push и pop соответственно с Varnode valueName.

emitAssignVarnodeFromVarnode (String varnodeOutName, String varnodeInName, int size)

Генерирует p-code для операции присвоения значения varnodeOutName = varnodeInName

emitVarnodeCall (String target, int size)

Генерирует P-Code для вызова функции target.

Заключение

Благодаря вышеизложенному механизму у нас получилось значительно улучшить вывод декомплилятора Ghidra. В итоге динамическая генерация p-code стала еще одним кирпичиком в нашем большом инструменте модуле для анализа скомпилированного bytenode скриптов Node.JS. Исходный код модуля доступен в нашем репозитории на github.com. Пользуйтесь, и удачного вам реверс-инжиниринга!

Если у вас остались какие-то вопросы, задавайте их в комментариях - буду рад ответить.

Большое спасибо за исследование особенностей Node.js и разработку модуля моим коллегам: Владимиру Кононовичу, Наталье Тляповой, Сергею Федонину.

Подробнее..

Получение исходного кода PowerPacker Cruncher от AmigaOS

10.08.2020 22:11:31 | Автор: admin


Всем привет,


Демо-сцена существует очень давно. Зачастую, в процессе разработки очередной крутой демки приходится изобретать крутые алгоритмы: как для красивых анимаций и трекерной музыки, так и для кода. Иногда код получается большого объёма, поэтому его требуется сжать.


Понятно, что можно взять любой доступный алгоритм сжатия и использовать его у себя, но не существовало бы сейчас такого огромного количества различных упаковщиков, если бы всем хватало одного единственного алгоритма. Кому-то не нравится скорость работы, кому-то качество сжатия, вот и изобретаются всё новые и новые алгоритмы. Одним из них и стал PowerPacker, исходные коды которого хотели получить многие, но удалось только мне.


Немного о PowerPacker


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


Сначала PowerPacker распространялся в виде самостоятельных исполняемых файлов: упаковщика и распаковщика. Затем, похоже, спрос на данный алгоритм сжатия вырос, и автор (Nico Franois) решил сделать своё творение платным, при этом перейдя на распространение уже в виде библиотеки powerpacker.library.


Получение исходников


Для получения исходников, как и в случае с RNC ProPack, пришлось написать множество вспомогательного инструментария:


  1. Плагин-отладчик для IDA Pro (не работает, забросил)
  2. Загрузчик Amiga Hunk для Ghidra (помог)
  3. Загрузчик для library-файлов для Ghidra (очень помог)
  4. gdb-сервер для AmigaOS, работающий на ней же (не работал на моих файлах)

Отдельным пунктом идёт покупка kickstart rom (это что-то типа биоса, нутрянки AmigaOS, без него работать ничего не будет).


Потом у IDA появилась возможность отлаживать через GDB в том числе и для m68k. Правда серверной части, которая могла бы при этом эмулировать и мои файлы, и AmigaOS, у меня не было. WinUAE не умеет в gdb до сих пор.


Затем, спустя несколько лет, появилось расширение для Visual Code: https://github.com/BartmanAbyss/vscode-amiga-debug, которое позволяет отлаживать исходные файлы на C, при помощи модифицированного WinUAE с добавленным в него gdb-сервером. Вот здесь я и осознал шанс на декомпиляцию есть.


Декомпиляция


Этот процесс без собственно самого декомпилятора превращается в долгое и мучительное преобразование ассемблерных инструкций в сишный код. И, если с кодом, который генерировался C-компилятором, проблем обычно не возникает, то вот с вручную написанным ассемблерным кодом проблем достаточно. Вот самые основные из них:


  • циклы (бесконечные goto)
  • использование одного и того же регистра как для хранения 16-битных значений, так и для хранения 32-битных. А ещё они в какой-то момент становятся знаковыми, хотя до этого использовались как беззнаковые.

Отладочный стенд


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


C:\Users\<USER>\.vscode\extensions\bartmanabyss.amiga-debug-1.0.0

Создаём и компилируем тестовый пример (да, у расширения он имеется). В подпапке .\bin имеется следующий список файлов:


  • dh0\
  • dh0\runme.exe
  • dh0\s\
  • dh0\s\startup-sequence
  • opt\
  • default.uae
  • elf2hunk.c
  • elf2hunk.exe
  • gnumake.exe
  • winuae.ini
  • winuae-gdb.exe

Подкаталог .\dh0\s содержит файл startup-sequence, в котором указываются команды, запускаемые при старте операционной системы. У меня он выглядит вот так:


:runme.exe

Здесь можно добавить нужные аргументы или команды. Для моих целей необходимо заменить файл runme.exe на исполняемый файл от PowerPacker-а, который затем будет загружать ту самую powerpacker.library. А вот куда класть эту библиотеку я понял не сразу. Оказывается, нужно было создать в каталоге .\dh0\ подкаталог Libs (я подсмотрел эту структуру в уже запущенной AmigaOS) и положить туда. Запускаю.



После выполнения данной команды произойдёт запуск winuae-gdb.exe, открытие порта 2345 для работы с gdb, и остановка на точке входа запускаемой программы. Остаётся только подключиться с помощью IDA и её Remote GDB debugger к сессии WinUAE.



Меняем порт на 2345, жмём Debugger->Attach to process..., затем выбираем процесс с id = 0.



После этого у нас появляется окно отладки:



Как видим, адрес на котором мы стоим, отличается от адреса, на котором создавалась idb 0x10000, поэтому останавливаем отладку и делаем Rebase на 0x27D30. Это поможет в дальнейшей отладке не терять изменений, сделанных в базе.


С этого момента можно спокойно заниматься пошаговой отладкой до тех пор, пока вы не превысите количество брейкпоинтов равное 20. Сначала я не догадывался, в чём причина, но мои брейкопоинты вдруг становились неактивными, невалидными. Лишь посмотрев в исходник WinUAE (который, к тому же, собрать совершенно не просто), я нашёл ограничение в 20 брейкопоинтов. Собрав новую сборку с количеством, равным 999, мне удалось наконец-то безболезненно заниматься самим процессом отладки.


Библиотека powerpacker.library


Тут пришлось изощряться, попутно найдя изящное решение, которое может помочь и вам при отладке загружаемых библиотек. Дело в том, что загруженные в память библиотеки (как и другие появляющиеся только во время отладки регионы памяти), можно сохранять прямо в idb, и работать с ними, при желании, в статике. При этом, при перезапуске процесса отладки, вы не потеряете свои наработки по переименованию переменных, меток, и т.п. Для проворачивания этой хитрости необходимо на необходимом сегменте, после загрузки нужной библиотеки, зайти в его свойства (выбрав Edit segment...):




Вы увидите, что там присутствует галка Debugger segment, при снятии которой и нажатии OK, данный сегмент будет сохранён в базу. Единственный момент: стоит следить за размером сегмента, иначе сохранение его в базу может растянуться, или вообще не закончиться.


Теперь можно входить в вызовы экспортируемых функций, и, при одном и том же адресе загрузки библиотеки, вы будете попадать в свой, уже проанализированный, код. Удобно.


В случае AmigaOS вызовы экспортируемых функций выглядят как вызовы по отрицательным смещениям относительно базы, по которой загружена библиотека:



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


Результаты работы


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


Протестировав всё на 210 файлах, найдя и исправив другие вылезающие баги (такие как выход за границы массива в оригинальном алгоритме), я готов опубликовать результаты своей работы:



Ссылки


Исходники: https://github.com/lab313ru/powerpacker_src
Релизы: https://github.com/lab313ru/powerpacker_src/releases

Подробнее..

Из песочницы Пишем спецификацию под Nvidia Kepler (бинарники CUDA, версия языка sm_30) для Ghidra

02.07.2020 00:11:00 | Автор: admin
Для обычных процессорных языков уже написано довольно много спецификаций для Ghidra, однако для графических ничего нет. Оно и понятно, ведь там своя специфика: предикаты, константы, через которые передаются параметры в том числе, и другие вещи, унаследованные от шейдеров. Кроме того формат, который используется для хранения кода, зачастую проприетарный, и его нужно самостоятельно ревёрсить.

В этой статье на двух примерах разберёмся, что к чему.

Первая программа простейший 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
Подробнее..

Категории

Последние комментарии

  • Имя: Макс
    24.08.2022 | 11:28
    Я разраб в IT компании, работаю на арбитражную команду. Мы работаем с приламы и сайтами, при работе замечаются постоянные баны и лаги. Пацаны посоветовали сервис по анализу исходного кода,https://app Подробнее..
  • Имя: 9055410337
    20.08.2022 | 17:41
    поможем пишите в телеграм Подробнее..
  • Имя: sabbat
    17.08.2022 | 20:42
    Охренеть.. это просто шикарная статья, феноменально круто. Большое спасибо за разбор! Надеюсь как-нибудь с тобой связаться для обсуждений чего-либо) Подробнее..
  • Имя: Мария
    09.08.2022 | 14:44
    Добрый день. Если обладаете такой информацией, то подскажите, пожалуйста, где можно найти много-много материала по Yggdrasil и его уязвимостях для написания диплома? Благодарю. Подробнее..
© 2006-2024, personeltest.ru