使用 GPU 的程序（程序集）是什么样子的？

Question

从这个答案看来，GPU制造商只是为特定的GPU API提供了驱动程序，并且不存在GPU组装这样的东西，或者至少，永远不会有像AMD64程序员手册那样出版的GPU组装编程手册

然而，据我了解，所有进程运行都经过CPU，并且可以被反汇编。

我的问题是：使用 GPU 的程序汇编会是什么样子？我的假设是，它将使用系统调用来操作代表 GPU 的设备文件。这个假设正确吗？

Answer 1

使用 GPU 的代码是什么样的？

阅读有关OpenCL的更多信息（或者，仅对于 Nvidia 硬件，了解有关CUDA 的信息）。还要注意OpenACC！另请参阅OpenCL 相关资源，并阅读一些 OpenCL 书籍。阅读一些OpenCL 教程。

实际上，您永远不会看到GPGPU的“汇编代码” 。但您将使用 OpenCL 进行编码（它的级别非常低，并且根据特定硬件调整代码非常困难且容易出错）。

AFAIK，AMD倾向于发布其大多数GPU的“机器代码规范”（例如ISA ）。英伟达则要神秘得多。请注意，SPIR是“类似汇编的”（实际上基于LLVM 字节码），但仍然不完全是汇编程序。

我的问题是：使用 GPU 的程序汇编会是什么样子？我的假设是，它将使用系统调用来操作代表 GPU 的设备文件。这个假设正确吗？

系统调用（非常特定于硬件）将 SPIR 或等效字节码（通常是 GPGPU 特定的机器代码）从 CPU（和虚拟内存）传输到 GPU，并将数据从GPGPU 传输到 CPU（和内存）并返回。细节非常复杂，而且通常是硬件制造商专有的。您更喜欢使用 OpenCL（或 CUDA）API 和方言。你的假设是错误的，或者至少过于简单化到毫无意义的地步。

另请参阅osdev.org wiki。

实际上，一些开源数值库（例如TensorFlow、OpenCV、BLAS等）都有 OpenCL 后端。因此需要花费几个月的时间来研究他们的源代码。

了解所有细节将为您提供博士学位。艾伯特·科恩（Albert Cohen ）（和许多其他专家）可能是您的顾问。

另请阅读有关AMDGPU及其GCN的更多信息。例如，查看AMD Vega规范。

然而，据我了解，所有进程运行都经过CPU，并且可以被反汇编。

这是一个非常天真的说法，我认为这是错误的（至少对于我喜欢编写的程序来说，它们都以某种方式在运行时生成代码）。在实践中，你不会理解反汇编的代码（这就是反编译如此困难的原因）。例如，生成机器代码的程序，请查看（在 Linux 上）SBCL（其REPL在每次用户交互时发出机器代码），或任何元程序，或大多数使用JIT 编译技术的程序（实际上，大多数 Java JVM正在进行 JIT 翻译）。我的manydl.cLinux 程序在运行时生成 C 代码，将其编译成共享库，即可以动态链接的插件，然后dlopen(3) -ing 该插件（并且可以重复数十万次）。有关有助于生成机器代码的库的示例，请参阅libgccjit。

您还应该阅读有关操作系统的更多信息。我强烈推荐操作系统：三个简单的部分（可免费下载）。

Answer 2

如果您使用的是 Nvidia GPU，那么您可以查看PTX汇编代码。PTX 只是伪汇编，介于 OpenCL 和 GPU 上实际运行的二进制代码之间。这是从 OpenCL 获取它的方式：

Context context(device);
queue = CommandQueue(context, device); // queue to push commands for the device
Program::Sources source;
string kernel_code = opencl_code_settings(N,M)+opencl_code();
source.push_back({ kernel_code.c_str(), kernel_code.length() });
Program program(context, source);
if(program.build("-cl-fast-relaxed-math")) return false; // compile OpenCL code, return false if there is an error
const string ptx_code = program.getInfo<CL_PROGRAM_BINARIES>()[0]; // generate assembly (ptx) for OpenCL code

该字符串ptx_code就是您要查找的内容。这是一个小内核示例：

kernel void benchmark_1(global float* data) {
    const uint n = get_global_id(0);
    #pragma unroll
    for(uint i=0; i<def_M; i++) data[i*def_N+n] = 0.0f;
}

该内核的 PTX 代码如下所示：

//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: UNKNOWN
// Driver 
// Based on LLVM 3.4svn
//

.version 6.2
.target sm_61, texmode_independent
.address_size 64

    // .globl   benchmark_1

.entry benchmark_1(
    .param .u64 .ptr .global .align 4 benchmark_1_param_0
)
{
    .reg .b32   %r<23>;
    .reg .b64   %rd<34>;


    ld.param.u64    %rd1, [benchmark_1_param_0];
    mov.b32 %r1, %envreg3;
    mov.u32     %r2, %ntid.x;
    mov.u32     %r3, %ctaid.x;
    mad.lo.s32  %r4, %r3, %r2, %r1;
    mov.u32     %r5, %tid.x;
    add.s32     %r6, %r4, %r5;
    mul.wide.u32    %rd2, %r6, 4;
    add.s64     %rd3, %rd1, %rd2;
    mov.u32     %r7, 0;
    st.global.u32   [%rd3], %r7;
    add.s32     %r8, %r6, 15728640;
    mul.wide.u32    %rd4, %r8, 4;
    add.s64     %rd5, %rd1, %rd4;
    st.global.u32   [%rd5], %r7;
    add.s32     %r9, %r6, 31457280;
    mul.wide.u32    %rd6, %r9, 4;
    add.s64     %rd7, %rd1, %rd6;
    st.global.u32   [%rd7], %r7;
    add.s32     %r10, %r6, 47185920;
    mul.wide.u32    %rd8, %r10, 4;
    add.s64     %rd9, %rd1, %rd8;
    st.global.u32   [%rd9], %r7;
    add.s32     %r11, %r6, 62914560;
    mul.wide.u32    %rd10, %r11, 4;
    add.s64     %rd11, %rd1, %rd10;
    st.global.u32   [%rd11], %r7;
    add.s32     %r12, %r6, 78643200;
    mul.wide.u32    %rd12, %r12, 4;
    add.s64     %rd13, %rd1, %rd12;
    st.global.u32   [%rd13], %r7;
    add.s32     %r13, %r6, 94371840;
    mul.wide.u32    %rd14, %r13, 4;
    add.s64     %rd15, %rd1, %rd14;
    st.global.u32   [%rd15], %r7;
    add.s32     %r14, %r6, 110100480;
    mul.wide.u32    %rd16, %r14, 4;
    add.s64     %rd17, %rd1, %rd16;
    st.global.u32   [%rd17], %r7;
    add.s32     %r15, %r6, 125829120;
    mul.wide.u32    %rd18, %r15, 4;
    add.s64     %rd19, %rd1, %rd18;
    st.global.u32   [%rd19], %r7;
    add.s32     %r16, %r6, 141557760;
    mul.wide.u32    %rd20, %r16, 4;
    add.s64     %rd21, %rd1, %rd20;
    st.global.u32   [%rd21], %r7;
    add.s32     %r17, %r6, 157286400;
    mul.wide.u32    %rd22, %r17, 4;
    add.s64     %rd23, %rd1, %rd22;
    st.global.u32   [%rd23], %r7;
    add.s32     %r18, %r6, 173015040;
    mul.wide.u32    %rd24, %r18, 4;
    add.s64     %rd25, %rd1, %rd24;
    st.global.u32   [%rd25], %r7;
    add.s32     %r19, %r6, 188743680;
    mul.wide.u32    %rd26, %r19, 4;
    add.s64     %rd27, %rd1, %rd26;
    st.global.u32   [%rd27], %r7;
    add.s32     %r20, %r6, 204472320;
    mul.wide.u32    %rd28, %r20, 4;
    add.s64     %rd29, %rd1, %rd28;
    st.global.u32   [%rd29], %r7;
    add.s32     %r21, %r6, 220200960;
    mul.wide.u32    %rd30, %r21, 4;
    add.s64     %rd31, %rd1, %rd30;
    st.global.u32   [%rd31], %r7;
    add.s32     %r22, %r6, 235929600;
    mul.wide.u32    %rd32, %r22, 4;
    add.s64     %rd33, %rd1, %rd32;
    st.global.u32   [%rd33], %r7;
    ret;
}

例如，您可以通过 PTX 代码对FLOP 和内存传输进行计数，以检查代码通过 Roofline 模型运行的效率。