如何创建或操作 GPU 汇编器？答案

【问题标题】：How to create or manipulate GPU assembler?如何创建或操作 GPU 汇编器？
【发布时间】：2011-06-07 08:29:35
【问题描述】：

是否有人在创建/操作 GPU 机器代码（可能是在运行时）方面有经验？

我对修改 GPU 汇编代码很感兴趣，可能在运行时以最小的开销进行。具体来说，我对基于汇编程序的基因编程很感兴趣。

我知道 ATI 已经为他们的一些卡发布了 ISA，并且 nvidia 最近发布了针对旧卡的 CUDA 反汇编程序，但我不确定是否可以在运行时甚至之前修改内存中的指令。

这可能吗？欢迎提供任何相关信息。

【问题讨论】：

你有nvidia最近发布的反汇编器的链接吗？我发现的只有“decuda”，这是一部独立的作品；我认为 nvidia 从未发布过有关其硬件实际理解的操作码的信息。
它可能只发布给注册开发者，虽然我认为他们包含在最新的 CUDA 版本中
cuobjdump 只允许您从主机二进制文件中提取 *.cubin 文件或线性反汇编。 FERMI 没有完整的参考，例如 x86。或者你能告诉我当我们执行减法时设置了哪些标志？

标签： assembly stream cuda gpgpu

【解决方案1】：

在 CUDA 驱动程序 API 中，module management functions 允许应用程序在运行时加载“模块”，它（大致）是 PTX 或 cubin 文件。 PTX 是中间语言，而 cubin 是已经编译好的指令集。 cuModuleLoadData() 和 cuModuleLoadDataEx() 似乎能够从 RAM 中的指针“加载”模块，这意味着不需要实际的文件。

所以您的问题似乎是：如何以编程方式在 RAM 中构建 cubin 模块？据我所知，NVIDIA 从来没有公布过他们的硬件能够真正理解的指令的细节。但是，有一个名为decuda 的独立开源包，其中包括“cudasm”，这是“旧”NVIDIA GPU 理解的汇编程序（“旧”= GeForce 8xxx 和 9xxx）。我不知道集成到更广泛的应用程序中会有多容易。它是用 Python 编写的。

较新的 NVIDIA GPU 使用不同的指令集（有多少不同，我不知道），因此旧 GPU 的 cubin（NVIDIA/CUDA 术语中的“计算能力 1.x”）可能无法在最近的 GPU 上运行（计算能力 2.x，即“Fermi 架构”，例如 GTX 480）。这就是为什么通常首选 PTX 的原因：给定的 PTX 文件可以跨 GPU 代移植。

【讨论】：

【解决方案2】：

NVIDIA Fermi ISA 的汇编程序：http://code.google.com/p/asfermi

【讨论】：

【解决方案3】：

我发现gpuocelot 开源（BSD 许可证）项目很有趣。

它是“PTX 的动态编译框架”。我称它为 cpu 翻译器。

“Ocelot 目前允许在 NVIDIA GPU、AMD GPU 和 x86-CPU 上执行 CUDA 程序”。据我所知，该框架对 PTX Kernel 进行控制流和数据流分析，以便应用适当的转换。

【讨论】：

【解决方案4】：

NVIDIA PTX 生成和修改

不确定它与硬件相比有多低（可能没有记录？），但它可以从类似 C/C++ 的 GPU 语言生成，并以几种方式修改和重用：

OpenCL clGetProgramInfo(program, CL_PROGRAM_BINARIES + clCreateProgramWithBinary：最小可运行示例：How to use clCreateProgramWithBinary in OpenCL?

这些是标准化的 OpenCL API，它们产生和使用实现定义的格式，在 Linux 驱动程序版本 375.39 中恰好是人类可读的 PTX。

因此您可以转储 PTX、修改它并重新加载。
nvcc: 可以简单地将 CUDA GPU 端代码编译为 ptx 程序集：
```
nvcc --ptx a.cu
```
nvcc 也可以编译包含设备和主机代码的 OpenCL C 程序：Compile and build .cl file using NVIDIA's nvcc Compiler? 但我找不到如何使用 nvcc 获取 ptx。哪种有意义，因为它只是普通的 C + C 字符串，而不是神奇的 C 超集。这也是由：https://arrayfire.com/generating-ptx-files-from-opencl-code/
提出的
而且我不确定如何重新编译修改后的 PTX 并像使用 clCreateProgramWithBinary 一样使用它：How to compile PTX code

使用clGetProgramInfo，一个输入CL内核：

__kernel void kmain(__global int *out) {
    out[get_global_id(0)]++;
}

被编译成一些 PTX，例如：

//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-21124049
// Cuda compilation tools, release 8.0, V8.0.44
// Based on LLVM 3.4svn
//

.version 5.0
.target sm_20
.address_size 64

    // .globl   _Z3incPi

.visible .entry _Z3incPi(
    .param .u64 _Z3incPi_param_0
)
{
    .reg .pred  %p<2>;
    .reg .b32   %r<4>;
    .reg .b64   %rd<5>;


    ld.param.u64    %rd1, [_Z3incPi_param_0];
    mov.u32     %r1, %ctaid.x;
    setp.gt.s32 %p1, %r1, 2;
    @%p1 bra    BB0_2;

    cvta.to.global.u64  %rd2, %rd1;
    mul.wide.s32    %rd3, %r1, 4;
    add.s64     %rd4, %rd2, %rd3;
    ldu.global.u32  %r2, [%rd4];
    add.s32     %r3, %r2, 1;
    st.global.u32   [%rd4], %r3;

BB0_2:
    ret;
}

然后，例如，如果您修改该行：

add.s32     %r3, %r2, 1;

到：

add.s32     %r3, %r2, 2;

并重用修改后的 PTX，它实际上增加了 2 而不是预期的 1。

【讨论】：

@talonmies 感谢您的反馈！编译 OpenCL 需要实际的 C 程序，就像 CUDA 一样。请参阅：stackoverflow.com/questions/13062469/… 但是我说你可以为 OpenCL 提取 ptx 和 nvcc 是错误的，它只适用于 CUDA（我同时测试了太多东西）。 clGetProgramInfo 一直在工作，正如所述。我已经更新了更清楚地解释这些观点的答案，并取消了它。如果您发现有任何问题，请告诉我。

【解决方案5】：

OpenCL 就是为此目的而设计的。您将程序作为字符串提供，并可能在运行时对其进行编译。请参阅其他海报提供的链接。

【讨论】：

据我所知，OpenCL 在安装时首先编译成中间语言 IL（类似于 NVidia 的 PTX），然后正确编译成机器指令。这是我感兴趣的机器指令。
不，您可以像我写的那样从字符串中即时编译 OpenCL。