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

Sleigh

Создаем процессорный модуль под Ghidra на примере байткода v8

23.04.2021 18:05:42 | Автор: admin

В прошлом году наша команда столкнулась с необходимостью анализа байткода V8. Тогда еще не существовало готовых инструментов, позволявших восстановить такой код и обеспечить удобную навигацию по нему. Было принято решение попробовать написать процессорный модуль под фреймворк Ghidra. Благодаря особенностям используемого языка описания инструкций на выходе мы получили не только читаемый набор инструкций, но и C-подобный декомпилятор. Эта статья продолжение серии материалов (1, 2) о нашем плагине для Ghidra.

Между написанием процессорного модуля и статьи прошло несколько месяцев. За это время спецификация SLEIGH не изменилась, и описанный модуль работает на версиях 9.1.29.2.2, которые были выпущены за последние полгода.

Сейчас на ghidra.re и в приложенной к Ghidra документации есть достаточно хорошее описание возможностей языка эти материалы стоит почитать перед написанием своих модулей. Отличными примерами могут быть уже готовые процессорные модули разработчиков фреймворка, особенно если вы знаете описываемую в них архитектуру.

В документации можно прочесть, что процессорные модули для Ghidra пишутся на языке SLEIGH, который произошел от языка SLED (Specification Language for Encoding and Decoding) и разрабатывался целенаправленно под Ghidra. Он транслирует машинный код в p-code (промежуточный язык, используемый Ghidra для построения декомпилированного кода). Как у языка, предназначенного для описания инструкций процессора, у него достаточно много ограничений, которые, однако, можно купировать за счет механизма внедрения p-code в java-коде.

Исходный код созданного процессорного модуля представлен на github. В этой статье будут рассматриваться принципы и ключевые понятия, которые использовались при разработке процессорного модуля на чистом SLEIGH на примере некоторых инструкций. Работа с пулом констант, инъекции p-code, анализатор и загрузчик будут или были рассмотрены в других статьях. Также про анализаторы и загрузчики можно почитать в книге The Ghidra Book: The Definitive Guide.

С чего начать

Для работы понадобится установленная среда разработки Eclipse, в которую нужно добавить плагины, поставляемые с Ghidra: GhidraDev и GhidraSleighEditor. Далее создается Ghidra Module Project с именем v8_bytecode. Созданный проект содержит шаблоны важных для процессорного модуля файлов, которые мы будем модифицировать под свои нужды.

Чтобы получить общее представление о файлах, с которыми предстоит работать, обратимся к официальной документации либо вышедшей недавно книге Криса Игла и Кары Нанс The Ghidra Book: The Definitive Guide. Вот описание этих файлов.

  • *.сspec спецификация компилятора.

  • *.ldefs определение языка. Содержит отображаемые в интерфейсе параметры процессорного модуля. Также содержит ссылки на файлы *.sla, спецификацию процессора и спецификации компилятора.

  • *.pspec спецификация процессора.

  • *.opinion конфигурации для загрузчика; поскольку мы будем описывать только один вид файлов, файл opinion можно оставить пустым: он не пригодится.

  • *.slaspec, *.sinc файлы, описывающие регистры и инструкции процессора на языке SLEIGH.

Также после первого запуска вашего проекта появится файл с расширением .sla, он генерируется на основании slaspec-файла.

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

О регистрах V8

Jsc-файл, который нас интересовал, был собран c использованием среды выполнения JavaScript Node.Js 8.16.0 через bytenode (этот модуль либо будет в поставке Node.Js, либо нужно будет доставить его через npm). По сути, bytenode использует документированный функционал Node.js для создания скомпилированного файла. Вот исходный код функции, компилирующей jsc файлы из js:

Node.js можно скачать как в собранном виде, так и в виде исходников. При детальном изучении исходных файлов и примеров инструкций становится ясно, как кодируются регистры в байткоде (для понимания расчета индекса будут полезны файлы bytecode-register.cc, bytecode-register.h). Примеры инструкций v8 с расчетами индексов регистров в соответствии с Node.js:

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

Тут Х количество аргументов текущей функции без учета передаваемого <this>, aX регистры, содержащие аргументы функции, а rN регистры, используемые как локальные переменные. Регистры могут кодироваться 1-байтовыми значениями для обычных инструкций, 2-байтовыми для инструкций с пометкой Wide- и 4-байтовыми для инструкций с пометкой ExtraWide-. Пример кодировки Wide-инструкции с пояснениями:

Более подробно о Node.js и v8 можно почитать в статье Сергея Федонина.

Стоит заметить, что SLEIGH не совсем подходит для описания подобных интерпретируемых байткодов, поэтому у написанного процессорного модуля есть некоторые ограничения. Например, определена работа не более чем с 124регистрами rN и 125регистрами aX. Была попытка решить эту проблему через стековую модель взаимодействия с регистрами, так как она больше соответствовала концепции. Однако в этом случае дизассемблированный байткод тяжелее читался:

Также без введения дополнительных псевдоинструкций, регистров или областей памяти не представляется возможным высчитывать название регистра аргумента в соответствии с Node.js из-за отсутствия информации о количестве аргументов. В связи с этим нами было принято решение проставлять номера в названии регистров аргументов функций (X в aX) в обратном порядке. Это не мешает разбору кода, что было для нас важным критерием, однако может смущать при сравнении результатов вывода инструкций файла в разных инструментах.

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

CSPEC

Немного текстовой информации о тегах, используемых в cspec-файлах, можно найти в исходниках фреймворка на github. В описании назначения файла говорится:

Спецификация компилятора является необходимой частью модуля языка Ghidra для поддержки разборки и анализа конкретного процессора. Его цель закодировать информацию о целевом двоичном файле, специфичном для компилятора, сгенерировавшего этот двоичный файл. В Ghidra спецификация SLEIGH позволяет декодировать машинные инструкции для конкретного процессора, например Intelx86, но эти инструкции могут продуцировать более одного компилятора. Для конкретного целевого двоичного файла понимание деталей о конкретном компиляторе, используемом для его сборки, важно для процесса разбора кода. Спецификация компилятора удовлетворяет эту потребность, позволяя формально описывать такие концепции, как соглашения о передаче параметров и механизмы стека.

Также становится понятно, что теги используются для следующих целей:

  • Compiler Specific P-code Interpretation;

  • Compiler Datatype Organization (у нас использовался <data_organization>);

  • Compiler Scoping and Memory Access (у нас использовался <global>);

  • Compiler Special Purpose Registers (у нас использовался <stackpointer>);

  • Parameter Passing (у нас использовался <default_proto>).

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

Теги <data_organization> и <stackpointer> достаточно типовые; разберем тег <prototype> в <default_proto>, частично описывающий соглашение о вызове функций. Для него определим: <input>, <output>, <unaffected>.

Как говорилось выше, аргументы в функцию передаются через регистры aX. В модуле регистры должны быть определены как непрерывная последовательность байтов по смещению в некотором пространстве. Как правило, в таких случаях используется специально придуманное для этого пространство register. Однако теоретически не запрещено использовать любое другое. В случае наличия большого количества регистров, выполняющих примерно одни функции, проще всего не прописывать каждый отдельно, а просто указать смещение в пространстве регистров, по которому они будут определены. Поэтому в спецификации компилятора помечаем область памяти в пространстве регистров (space="register") в теге <input> для регистров, через которые происходит передача аргументов в функции, по смещению 0x14000 (0x14000 не несет в себе сакрального смысла, это просто смещение, по которому в *.slaspec далее будут определены регистры aX).

По умолчанию результат вызова функций сохраняется в аккумулятор (acc), что нужно прописать в теге <output>. Для альтернативных вариантов регистров, в которые происходит сохранение возвращаемых функциями значений, можно определить логику при описании инструкций. Отметим в теге <unaffected>, что вызовы функций на регистр, хранящий указатель на стек, не влияют.

Для работы с частью регистров наиболее удобным будет вариант определения их как изменяемых глобально, поэтому в теге <global> определяем диапазон регистров в пространстве register по смещению 0x2000.

LDEFS

Перейдем к определению языка это файл с расширением .ldefs. Он требует немного информации для оформления: порядок байт (у нас le), названия ключевых файлов (*.sla, *.pspec,*.cspec), id и название байткода, которое будет отображаться в списке поддерживаемых процессорных модулей при импорте файла в Ghidra. Если когда-то понадобится добавить процессорный модуль для файла, скомпилированного версией Node.js, существенно отличающейся от текущей, то имеет смысл описать его тут же через создание еще одного тега <language>, как это сделано для описания семейств процессоров в *.ldefs модулей, поставляемых в рамках Ghidra.

Практическое применение информации, не касающейся определения файлов, будет видно при попытке импорта файла.

PSPEC

Сложнее в плане документации дела обстоят со спецификацией процессора (файл с расширением .pspec). В данном случае можно обратиться к готовым решениям в рамках самого фреймворка или к файлу processor_spec.rxg (вариант с полноценным разбором исходных кодов Ghidra мы не рассматривали). Чего-то более подробного на момент написания модуля не было. Вероятно, со временем разработчики опубликуют официальную документацию.

В текущем проекте на данный момент от спецификации процессора может понадобиться только программный счетчик, оставим этот тег из стандартного шаблона свежесозданного проекта (на самом деле можно оставить <processor_spec> пустым).

SLASPEC

Теперь можно приступить к непосредственному описанию инструкций на SLEIGH в файле с расширением .slaspec.

Базовые определения и макросы препроцессора

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

Адресные пространства, которые понадобятся для описания байткода (у нас создаются пространства с именами register и ram), определяются через define space, а регистры через define register. Значение offset в определении регистров не принципиально, главное, чтобы они находились по разным смещениям. Занимаемое регистрами количество байтов определяется параметром size. Стоит помнить, что определенная тут информация должна соответствовать обращениям к аналогичным абстракциям и величинам в рамках *.cspec и анализатора, если вы ссылаетесь на эти регистры.

Описание инструкций

В документации (https://ghidra.re/courses/languages/html/sleigh_constructors.html) можно прочитать, что определение инструкций происходит через таблицы, которые состоят из одного и более конструкторов и имеют имена идентификаторы символов семейства. Таблицы в SLEIGH по сути являются тем, что называется символами семейства, в статье мы не будем углубляться в определения символов, для этих целей проще прочитать Знакомство с символами. Конструкторы состоят из 5частей.

  1. Table Header (заголовок таблицы)

  2. Display Section (секция отображения)

  3. Bit Pattern Sections (секция битового шаблона)

  4. Disassembly Actions Section (секция действий при дизассемблировании инструкций)

  5. Semantics Actions Section (семантическая секция)

Пока что звучит страшно, опишем основной смысл.

  1. В Table Header либо помечается идентификатор, который можно использовать в других конструкторах, либо ничего нет (тогда это непосредственно описание инструкции).

  2. Display Section шаблон, показывающий как выводить инструкцию в листинг Ghidra.

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

  4. Disassembly Actions Section дополняет секцию битового шаблона какими-то вычислениями, если ее в чистом виде недостаточно.

  5. Semantics Actions Section описывает, что делает эта инструкция по смыслу, чтобы показать это в декомпиляторе.

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

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

Несколько документированных особенностей секции отображения, которые понадобятся дальше:

  • ^ разделяет идентификаторы и/или символы в секции, между которыми не должно быть пробелов;

  • используются, чтобы вставлять жестко закодированные строки, которые не будут считаться идентификатором;

  • пробельные символы обрезаются в начале и конце секции, а их последовательности сжимаются в один пробел;

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

Токены и их поля

Для описания конструкторов инструкций необходимо определить битовые поля. Через них осуществляется привязка битов программы к определенным абстракциям языка, в который будет происходить трансляция. Такими абстракциями могут быть мнемоники, операнды и т.п. Определение полей происходит в рамках задания токенов, синтаксис их определения выглядит так:

Размер токена tokenMaxSize должен быть кратен8. Это может быть неудобно, если операнды или какие-то нюансы для инструкции кодируются меньшим количеством бит. С другой стороны, это компенсируется возможностью создавать поля разных размеров, кодирующих позиционно любые биты в пределах размеров, задаваемых токеном. Для таких полей должны соблюдаться условия: start- и endBitNumX находятся в диапазоне от 0 до tokenMaxSize-1 включительно и startBitNumX <= endBitNumX.

Для разбираемого байткода v8 не было необходимости создавать поля, отличные по размеру от токена. Но, если бы такие поля были и использовались совместно, они бы объединялись через логические операторы & или |.

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

Теперь опишем простейшую инструкцию байткода, не имеющую операндов. Определим поле, которое будет описывать опкод инструкции. Как видно выше в разделе про v8, код инструкции описывается одним байтом (есть также Wide- и ExtraWide- инструкции, но они здесь не будут рассматриваться, по сути они просто используют операнды больших размеров и дополнительные байты под опкод инструкции). Таким образом получаем:

Теперь, используя поле op для идентификации первого и единственного опкода, определяющего инструкции Illegal и Nop, пишем для них конструкторы:

Байт 0xa7 в листинге Ghidra отобразит как инструкцию Illegal, не имеющую операндов. Для этой инструкции в примере использовалось ключевое слово unimpl. Это неимплементированная команда, дальнейшая декомпиляция будет прервана, что удобно для отслеживания нереализованных семантических описаний. Для Nop оставлена пустая семантическая секция, то есть команда не повлияет на отображение в декомпиляторе, что и должна делать эта инструкция. На самом деле Nop не присутствует как инструкция в Node.js нашей версии, мы ввели ее искусственно для реализации функционала SwitchOnSmiNoFeedback, но об этом будет рассказано в статье Владимира Кононовича.

Описываем операнды и семантику

Усложним концепцию: опишем конструктор для операций LdaSmi, в рамках которой происходит загрузка целого числа в аккумулятор (acc в определении пространства регистров), и AddSmi, которая по сути представляет собой сложение значения в аккумуляторе c целым числом.

Для текущих и будущих нужд определим чуть больше полей на манер операндов в bytecodes.h Node.js, создадим их в новом токене с именем operand, поскольку у этих полей будет другое назначение. Создание нескольких полей с одинаковыми битовыми масками может быть обусловлено как удобством восприятия, так и использованием нескольких полей одного токена в рамках одной инструкции (см. пример с AddSmi).

С точки зрения листинга хочется видеть что-то наподобие LdaSmi [-0х2]. Поэтому определяем в секции отображения мнемонику, а в шаблон прописываем имена полей, которые должны подставляться из секции disassembly action или битового шаблона (квадратные скобки тут не обязательны, это просто оформление).

Для инструкции AddSmi в секции битового шаблона, помимо поля op, устанавливающего ограничение на опкод, через ; появляются поля из токена operand. Они будут подставлены в секцию отображения в качестве операндов. Маппинг на реальные биты происходит в том порядке, в котором поля указаны в секции битового шаблона. В семантической секции, используя документированные операции, реализуем логику инструкций (то, что делал бы интерпретатор, выполняя эти инструкции).

Через ; могут также, например, идти регистры, контекстные переменные (о них поговорим позже), комбинации полей одного токена или полей с контекстными переменными.

Вот так выглядит окно листинга с описанными инструкциями со включенным полем PCode в меню изменение полей листинга Ghidra. Окно декомпилятора пока что не будет показательным из-за оптимизации кода, поэтому на данном этапе стоит ориентироваться только на промежуточный p-code.

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

Инструкции возврата значения из функции реализуются с помощью ключевого слова return, и, как уже упоминалось ранее, чаще всего возвращение значения при вызове функции происходит через аккумулятор:

Выводим регистры по битовым маскам

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

В разделе Базовые определения и макросы препроцессора регистры уже были объявлены, но для того, чтобы нужные регистры выбирались в зависимости от представленных в байткоде бит, необходимо привязать их список к соответствующим битовым маскам. Поле kReg имеет размер 8бит. Через конструкцию attach variables последовательно определяем каким битовым маскам от 0b до 11111111b вышеприведенные регистры будут соответствовать в рамках последующего использования полей из заданного списка (в нашем случае только kReg) в конструкторах. Например, в этом описании видно, что операнд, закодированный как 0xfb (11111011b), интерпретируется при описании его через kReg как регистр r0.

Теперь, когда за переменной kReg закреплены регистры, ее можно использовать в конструкторах:

Усложним конструкцию для соответствия конструктора более высокоуровневым описаниям инструкций из interpreter-generator.cc исходников Node.js. Вынесем поле kReg в отдельный конструктор, идентификатор таблицы которого в Table Header назовем src. В его семантической секции появляется новое ключевое слово export. Если не вдаваться в детали построения p-code, то по смыслу export определяет значение, которое должно быть подставлено в семантическую секцию конструктора вместо src. Вывод в Ghidra не изменится.

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

Переходы по адресам с goto

В байткоде встречаются операции условного и безусловного перехода по смещению относительно текущего адреса. Для перехода по адресу или метке в SLEIGH используется ключевое слово goto. Примечательно для определения то, что в секции битового шаблона используется поле kUImm, однако оно не используется в чистом виде. В секцию отображения выводится просчитанное в disassembly action секции значение через идентификатор rel. Величина inst_start предопределена для SLEIGH и содержит адрес текущей инструкции.

Компиляция SLEIGH проходит. Правда, в таком варианте (листинг ниже), судя по выводу, не получается создать варноду (это объекты, которыми манипулирует p-code инструкция), содержащую привязку к конкретному пространству.

Воспользуемся рекомендуемым разработчиками способом и вынесем часть определения через создание дополнительного конструктора с идентификатором dest. Конструкция *[ram]:4 rel не обозначает, что мы берем 4байта по адресу rel. По факту экспортируется адрес rel в пространстве ram. Оператор * в SLEIGH обозначает разыменование, но в данном конкретном случае относится к нюансу создания варнод (подробнее в Dynamic References).

Указание пространства [ram] может быть опущено (пример в комментарии), так как при определении мы указали его пространством по умолчанию. Как видно в инструкциях p-code, смещение было помечено как принадлежащее ram.

Чуть сложнее выглядит инструкция JumpIfFalse из-за использования условной конструкции. В SLEIGH она используется вместе с ключевым словом goto. Для большего соответствия концепциям js величина False ранее была определена как регистр, и можно заметить, что в pspec диапазон пространства регистров, к которому она привязана, помечен как глобальный. Благодаря этому в псевдокоде она отображается в соответствии с именованием регистра, а не численным значением.

В рассмотренных примерах переход осуществляется по рассчитываемому относительно inst_start адресу. Рассмотрим инструкцию TestGreaterThan, в которой происходит переход с помощью goto к метке (<true> в примере ниже) и inst_next. Переход к метке в принципе должен быть интуитивно понятным: если условие истинно, то далее должны выполняться инструкции, следующие за местом ее расположения. Метка действительна в только в пределах ее семантической секции.

Конструкция goto inst_next фактически завершает обработку текущей инструкции и передает управление на следующую. Стоит обратить внимание, что для выполнения знакового сравнения используется s>, см. документацию.

Несколько регистровых операндов

В предыдущих инструкциях использовалось не более одного регистра в качестве операнда. Так что сначала поговорим об использовании большего количества регистров для этих целей.

Описание однотипных операндов через конструкторы, имеющие разные идентификаторы таблиц (см. конструкторы в примере ниже), может иметь практическое применение для использования в рамках одного конструктора корневой таблицы. Такой вариант применим не только с точки зрения соответствия описанию инструкций v8, но и для преодоления возможных ошибок. Например, 4операнда инструкции CallProperty2 являются регистрами, идентично задаваемыми с точки зрения битовой маски. Попытка определить конструктор как :CallProperty2 kReg, kReg, kReg, kReg, [kIdx] вызовет ошибку в компиляторе Sleigh при попытке открыть файл с помощью процессорного модуля. Поэтому в нашем модуле использовались конструкторы для создания чего-то наподобие алиасов:

Стоит отметить, конечно, что решить эту проблему также можно было без определения новых конструкторов. Например, определив и прописав поля callable, receiver, arg1 и arg2 в рамках какого-либо токена с последующей их привязкой через attach к списку регистров:

Каждое из этих полей работало бы аналогично kReg в предыдущих примерах. Какой именно способ использовать вопрос эстетики.

Вызовы функций

В инструкции CallProperty2 также примечательно то, что она в семантической секции использует конструкцию call [callable];, которую мы не использовали до этого. В v8 аргументы функции хранятся в регистрах aX (как мы и пометили в cspec). Однако, с точки зрения байткода, помещения туда аргументов непосредственно перед вызовом функции не происходит (случаи, когда это происходит, можно посмотреть в sinc-файле, например для x86). Интерпретатор делает это самостоятельно, ведь у него есть вся необходимая информация. Но ее нет у Ghidra, поэтому в семантической секции мы пропишем помещение аргументов в регистры вручную. Однако нам необходимо будет восстановить значения задействованных регистров после вызова, так как в вызывающей функции эти регистры тоже могут хранить какие-то значения, необходимые для потока выполнения. Можно сохранить их через локальные переменные:

Можно также применять вариант с сохранением аргументов в памяти (в данном случае на стеке: sp не используется инструкциями, потому не повлияет на отображение в декомпиляторе) при использовании макросов на примере CallUndefinedReceiver1:

При написании подобных модулей вместе с загрузчиком и анализатором стоит следить, чтобы порядок передаваемых аргументов совпадал с порядком аргументов функции, определяемых при их создании в java-коде. Также стоит отметить, что для нашей ситуации было полезно сохранять не больше аргументов функции, чем их количество в вызывающей функции, но это сложно реализовать на SLEIGH. Вариант решения проблемы можно будет прочитать в статье Вячеслава Москвина про внедрение p-code инструкций.

Определяемые пользователем операции

Чтобы не терять при декомпиляции инструкции, в которых не планируется или нет возможности описывать семантический раздел, можно использовать определяемые пользователем операции. Стоит отметить, что в ситуации с acc не требуется явно указывать размер, поскольку размер регистра определен явно, а использовать его не полностью тут не нужно. Однако при передаче в подобные пользовательские операции, например, числовых констант придется явно указывать размер передаваемого значения (как в примере с CallVariadicCallOther в разделе О диапазонах регистров далее по тексту). Пользовательские операции определяются как define pcodeop OperationName и используются в семантическом разделе конструкторов в формате, напоминающем вызов функции во многих языках программирования.

Эти операции могут использоваться для внедрения p-code-инструкций в анализаторе: вызовы добавляются через тег callotherfixup в cspec-файл и прописывается логика в анализаторе.

Без переопределения в java-коде пользовательские операции в декомпиляторе выглядят так же, как они определены в семантическом разделе:

Тестируем модуль

Уже на этом этапе можно попробовать проверить работу процессорного модуля на практике. Скомпилируем через bytenode jsc-файл из небольшого примера на js:

Попробуем запустить написанный на основании статьи проект и импортировать полученный jsc-файл в Ghidra. Если что-то описано неправильно, Ghidra выдаст ошибку, а в логах eclipse будет локализирован номер строки, вызвавшей ошибку. При исправлении кода не нужно перестраивать проект: sleigh перекомпилируется при следующей попытке открыть файл.

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

Поскольку мы описывали только инструкции, файл не будет проанализирован и разобран по функциям. Код находится по смещению 0х1с0. НажмемD для преобразования байтов в инструкции по этому смещению иF, чтобы создать функцию. Вот так она будет выглядеть:

Более понятным становится вывод при использовании других средств (помимо SLEIGH), доступных разработчикам модулей. Например, при добавлении работы с пулом констант (для обращения к нему в SLEIGH зарезервировано ключевое слово cpool) появится возможность разрезолвить числовой идентификатор в команде LdaGlobal. Вот так в последней версии нашего проекта выглядит функция (для сравнения):

Разумеется, было бы приятнее видеть большее соответствие исходному коду, написанному на JavaScript, однако этого нельзя добиться, описывая инструкции только в файлах .slaspec (и .sinc). Чуть больший простор для воображения откроет статья, в которой будет описан механизм внедрения операций p-code, позволяющий при полном доступе ко всем ресурсам приложения манипулировать инструкциями, из которых собирается дерево p-code. Как раз на основании созданного дерева p-code результат декомпиляции выстраивается и отображается в интерфейсе.

О диапазонах регистров

В байткоде v8 для ряда инструкций реализована работа с диапазонами, парами и/или тройками регистров. С точки зрения кодирования диапазонов присутствует байткод первого регистра и количество используемых в диапазоне регистров. Для пар и троек указывается только начальный регистр, так как интерпретатору заранее известно, сколько регистров надо использовать для данной инструкции.

На основании описания понятно, что достаточно простым решением было бы отобразить при разборе инструкций первый регистр и их количество, то есть примерно так: ForInPrepare r9, r10!3. Чуть большим компромиссом в пользу читаемости было бы выводить первый и последний регистры диапазона, но, забегая вперед, можно сказать, что с точки зрения реализации уже это потребовало бы использования таблиц, состоящих из нескольких конструкторов.

Таблицы, содержащие несколько конструкторов

В рамках проекта для удобства восприятия было решено отображать в листинге весь список передаваемых регистров. Для секции отображения нет готового шаблона для выведения диапазонов регистров. Можно руководствоваться принципами, аналогичными использованным для процессорного модуля ARM: распечатыванием переменных через цепочку конструкторов (только сам принцип, реализация нам не подойдет из-за разности архитектур).

Тут наглядно видно тот самый случай, когда таблица с неким идентификатором состоит из нескольких конструкторов. По сути, это несколько конструкторов с одинаковыми идентификаторами в секции заголовка таблицы с разными условиями в секции битового шаблона. Выбираться при работе с идентификатором в других конструкторах будет наиболее подходящий по условию вариант. Например, если описан конструктор с неким условием и без условий, то при истинности условия будет выбран первый вариант, хотя второй формально ему не противоречит, поскольку не накладывает вообще никаких условий.

Как можно предположить, глядя на побайтовое описание инструкции CallProperty выше, для раскрытия диапазона необходимо распечатать регистры, отталкиваясь от первого вхождения, ориентируясь на известный первый регистр диапазона и количество элементов в нем. С точки зрения секции отображения, диапазон создается из двух конструкторов: rangeSrc и rangeDst. rangeSrc своего рода инициализация, где мы сохраняем входные данные, rangeDst будет распечатывать регистры на основании полученной информации. И как раз для rangeDst понадобится создавать таблицы, содержащие несколько конструкторов: как минимум для отображения диапазонов на регистрах aX и rX отдельно.

Для реализации условий необходимо учесть ряд ограничений. Проверять значения в секции битового шаблона рекомендуется только через =, а уточнять тут значение напрямую регистра нельзя, как и присваивать ему значения в секции disassembly action. Это лишает нас возможности использовать какой-то временный регистр. Стартовый регистр и длина диапазона могут быть любыми, а реализоваться, как уже упоминалось, диапазон может как на регистрах aX, так и на rX, а также быть нулевой длины. Уже на этом этапе понятно: если мы не хотим создавать гигантское количество определений на все случаи жизни, было бы неплохо иметь некие счетчики, чтобы выяснить, сколько регистров выводить и с какой позиции.

Контекстные переменные

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

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

Важно отметить, что контекстные переменные должны объявляться до конструкторов.

В документации поясняется, что контекстные переменные, как правило, используются в секции битовых шаблонов для проверки наличия какого-то контекста и изменяются в секции disassembly action. Так что в конструкторе с идентификатором таблицы rangeSrc, которую мы будем использовать для отображения диапазонов, в disassembly action секции сохраняем код первого регистра диапазона в контекстную переменную offStart, а их количество в counter. В секции отображения обозначаем начало диапазона открывающейся скобкой {.

Также стоит отметить, что в v8 не используется регистр range_size: он введен искусственно для хранения размера диапазона, чтобы было удобнее работать с этим значением в рамках семантической секции конструктора инструкции. Именно rangeSrc поставляет стартовый регистр и размер диапазона для семантической секции инструкции.

В рамках таблицы с идентификатором rangeDst описано 5конструкторов для следующих случаев.

  • Код стартового регистра диапазона соответствует a0 и счетчик counter равен 0 (пустой диапазон).

  • Код стартового регистра диапазона соответствует r0 и счетчик counter равен 0 (пустой диапазон).

  • Код регистра диапазона в offStart совпадает с a0, в disassembly action секции счетчик counter уменьшается, код регистра в offStart увеличивается, переход к конструктору rangedst1.

  • Код регистра диапазона в offStart совпадает с r0, в disassembly action секции счетчик counter и код регистра в offStart уменьшаются, переход к конструктору rangedst1.

  • Код стартового регистра диапазона пока не найден, переход к конструктору rangedst1(для этого случая в конструкторе нет никаких условий, как упоминалось выше, он будет браться, когда остальные конструкторы не подошли).

В третьем и четвертом случаях происходит вывод регистра в отображение. Последующие конструкторы rangeDstN, где N натуральное число, состоят из тех же вариантов, только для регистров aN/rN.

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

В примере ниже описаны только rangeDst, rangeDst1, rangeDst2, чтобы не загромождать статью. Для получения представления о виде подобных таблиц этого достаточно, полную версию можно посмотреть в исходниках проекта на github. По сути, при работе с rangeDst будет проходиться цепочка конструкторов по возрастанию индекса Х в rangeDstX, пока не встретится стартовый регистр, а затем цепочка конструкторов, соответствующая по длине размеру выводимого диапазона.

Конструкторы с закрывающими фигурными скобками выглядят похожими, можно попробовать их объединить. Вспоминаем про использование логических операторов & и |.

Конструктор для CallProperty в готовом проекте выглядит так:

Вот что получается в листинге:

Возможно, сейчас сбивает с толку, что в семантической секции используется пользовательская операция CallVariadicCallOther. В проекте на github она была переопределена в java-коде инструкциями p-code. Использование инъекции p-code вместо реализации через операцию call было обусловлено желанием видеть список передаваемых аргументов в декомпиляторе (согласно исходникам Node.js, первый регистр диапазона является приемником, а остальные передаваемыми аргументами). Используя только slaspec, добиться этого было бы, мягко говоря, тяжело:

Если есть желание попробовать повторить реализацию диапазонов самостоятельно, можно описать семантику как:

Затем по аналогии можно доопределить конструкторы rangeDstХ (понадобится до r7 включительно) и уже тогда попробовать посмотреть, как выглядит скомпилированный код console.log(1,2,3,4,5,6). Можно собрать его самостоятельно через bytenode или забрать готовый тут. Функция будет находиться по смещению 0x167, а сама инструкция на 0x18b.

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

Стоит отметить, что в нашем проекте мы вынесли все конструкторы rangeDst в отдельный файл, чтобы не загромождать файл с описанием инструкций (как и расширенные инструкции, работающие операндами размером 2 и 4байта):

Итог

Разработанный процессорный модуль удовлетворяет требованиям, которые мы к нему предъявляли: инструмент дает возможность просматривать инструкции байткода для заданных функций файла, который необходимо было разбирать в рамках проекта. В качестве бонуса мы получили декомпилятор, пусть не идеальный, но позволяющий быстрее ориентироваться в логике приложения. Однако, как и при работе со многими процессорными модулями, в данном случае лучше сверяться непосредственно с инструкциями, а не слепо ему верить. Можно также отметить, что при наличии времени и желания улучшить инструмент стоит реализовать хранение типов, продумать концепцию импортов и устранить проблему с обратным порядком аргументов функций. Надеемся, это руководство упростит ваше погружение в написание процессорных модулей на языке SLEIGH.

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

Автор: Наталья Тляпова

Полезные ссылки:

  1. https://ghidra.re/courses/languages/html/sleigh.html документация на SLEIGH.

  2. https://github.com/NationalSecurityAgency/ghidra/tree/master/Ghidra/Framework/SoftwareModeling/data/languages полезные файлы с описаниями *.cspec, *.pspec, *.opinion, *.ldefs.

  3. https://spinsel.dev/2020/06/17/ghidra-brainfuck-processor-1.htmlхорошая статья о реализации модуля для brainfuck в Ghidra.

  4. https://github.com/PositiveTechnologies/ghidra_nodejs репозиторий с полной версией процессорного модуля для Ghidra с загрузчиком и анализатором.

Подробнее..

Из песочницы Пишем спецификацию под 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