为Ghidra编写Nvidia Kepler（CUDA二进制，sm_30语言版本）规范

对于常见的处理器语言，已经为Ghidra编写了很多规范，而对于图形化的则没有。这是可以理解的，因为它具有自己的特性：谓词，用于传递参数的常量，包括从着色器继承的其他内容。另外，用于存储代码的格式通常是专有的，需要自己反转。



在本文中，我们将弄清楚两个示例的含义。



第一个程序是最简单的axpy（GPGPU的Hello World模拟）。第二个有助于了解条件的实现并在GPU上跳转，因为 那里一切都不同。



所有的Nvidia语言都使用小的endian编码，因此请立即以8块的相反顺序将十六进制编辑器中的字节复制到某个记事本中（例如Notepad ++）（此处的指令长度是恒定的）。然后，通过程序员计算器（Microsoft的计算器是合适的），我们将其转换为二进制代码。接下来，我们寻找匹配项，组成指令的掩码，然后组成操作数。为了解码和搜索掩码，使用了十六进制编辑器和cuobjdump反汇编程序，有时需要一个汇编程序，如AMDGPU中一样（因为反汇编程序在那里不可用，但这是另一篇文章的主题）。它是这样工作的：尝试顺序地将计算器中的所有可疑位取反，然后为字节获取一个新的十六进制值，如果存在，我们将它们替换为通过nvcc或汇编程序编译的二进制文件，但并非总是如此。然后通过cuobjdump检查。



我以以下格式传播源代码（主要使用C语言，不带加号和OOP以便与机器GPU代码进行更紧密的通信），然后立即disasm +字节，因为这样做更加方便，它们不需要互换。



将其复制到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];
}

范例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;
}

不难猜测，这些测试最初是为机器代码设计的，因此编译器没有什么要优化的。对于其他所有内容，您将必须手动取消优化。在复杂的示例中，这可能根本不可能，因此在这种情况下，您将不得不依赖反编译器和前端。



通常，规则是这样的-为了测试前端，我们以任何简单的（尽可能少的优化）最适合（再现错误）的示例为例。对于其余部分，反编译的代码将已经进行了优化（或仅通过重构以某种方式对其进行了纠正）。但就目前而言，主要任务至少是与机器代码具有相同功能的正确代码。这是“软件建模”。 “软件建模”本身不涉及重构，将C转换为C ++，还原类，甚至不涉及诸如模式识别之类的事情。



现在，我们正在寻找助记符，操作数和修饰符的模式。



为此，请比较可疑指令（或字符串，如果更方便调用）之间的位（二进制）。您还可以利用其他用户在其stackoverflow问题上发布的内容，例如“帮助理解二进制/ sass /机器代码”，使用教程（包括中文）和其他资源。因此，主操作号存储在位58-63中，但是还有其他位0-4（它们区分指令“ I2F”，“ ISETP”，“ MOV32I”），而不是0-2（对于忽略，是3）空指令中有4位，在规范中将其标记为“ UNK”。



对于寄存器和常数，您可以尝试使用反汇编程序来查找影响转储输出的所有位，例如扰流器下方布置的位。我设法找到的所有字段都在Github的规范中，文件CUDA.slaspec，节标记。



然后，您需要提供寄存器的地址，它们又位于Github上。这是必要的，因为在微观层面上，Sleigh将类型为“ register_space”的寄存器注册为空间中的全局变量，但是由于它们的空间未标记为“可推断的”（并且很可能无法标记），然后在反编译器中它们成为局部变量（最常见的是带有“ Var”前缀，但有时也有“ local”前缀）或参数（ param_”）。 SP从来没有派上用场，通常需要正式使用SP以确保反编译器正常工作。需要一台PC（类似于x86的IP）进行仿真。



然后是谓词寄存器，类似于标志，但是比“预定目的”（例如溢出，（不）等于零等）更多的“通用目的”。

然后，一个阻塞寄存器模拟一堆指令ISCADD .CC和IMAD.HI，因为在我的实现中，它们中的第一个会自行计算，第二个会进行计算，以避免将总和的一部分转移到高4个字节，因为这会搞反编译。但是随后您需要锁定下一个寄存器，直到IMAD.HI操作完成。类似的东西官方文档和反编译器预期输出之间的差异已经在同一Ghidra的SPU中。



然后是特殊的寄存器，到目前为止，这些寄存器都是通过cpool实现的。将来，我计划将它们替换为默认的为某些“可推断”空间定义的字符。这些是相同的threadIdx，blockIdx。



然后，将变量绑定到字段dest，par0，par1，par2，res。然后是子表，在这些子表之后（首先启动所有工作）是带有主要指令的主（根）表。



在这里，您必须严格遵循“助记符-操作数”的格式，但是，修饰符可以免除，但是修饰符必须附加到助记符或带有操作数的部分。不允许使用其他格式，即使相同的Hexagon DSP asm也必须适应此语法，但这并不是很困难。



最后阶段将以Pcode微编程语言编写指令的实现。在第一个示例中，我唯一要注意的是ISCADD .CC和IMAD.HI指令，其中的第一个使用指向寄存器的指针，并将它们作为8个字节而不是4个字节的指针解除引用。这是有意进行的，目的是为了更好地适应反编译器及其行为，尽管Nvidia文档中写有关于部分金额转移的内容。



对于第二个示例，最好选中反编译器设置中与“简化谓词”字样相对的复选框。关键是谓词是不同指令的一个相同条件，本质上仅是众所周知的“ SIMD”或其下一个等价物。那些。如果谓词位被置位，那么指令将连续执行。



您还需要养成立即为每条指令编写实现的习惯，而不仅仅是原型（操作数助记符），因为还有反编译器，仿真器和其他分析器。

但是总的来说，用Pcode编写实现比编写用于字节解码器的语法要容易得多。很快，事实证明，由于非常方便的中间语言，单一的Middleland（优化程序），2个后端（主要是C;作为替代方案-Java / C＃，更类似于后者，即Java），修复了来自x86的一些复杂指令的实现。 （因为goto会不时出现，但不会出现标记的中断）。

在以后的文章中，可能还会有托管语言的前端，例如DXBC，SPIR-V，它们将使用Java / C＃后端。但是到目前为止，只有机器代码在计划中。字节码需要一种特殊的方法。Ghidra帮助



项目 ：Pcode Sleigh