如何创建可执行文件来运行给定 PTX 文件中的内核？

Question

如何创建可执行文件来运行给定 PTX 文件中的内核？

The*_*hem 1 cuda build ptx

据我所知，你需要一个主机代码（对于CPU）和一个设备代码（对于GPU），没有它们你就无法在GPU上运行某些东西。

我正在学习 PTX ISA，但不知道如何在 Windows 上执行它。我是否需要 .cu 文件来运行它，或者是否有其他方法来运行它？

Answer 1

Rob*_*lla 5

长话短说：

如何组装 .ptx 文件和主机代码文件并生成可执行文件？

您使用CUDA 驱动程序 API。相关示例代码是vectorAddDrv（或者可能是任何其他驱动程序 API 示例代码）以及ptxjit.

我是否需要 .cu 文件来运行它，或者是否有其他方法来运行它？

如果您从 PTX 形式的设备代码开始，则不需要.cu文件（也不需要nvcc）来使用驱动程序 API 方法。

细节：

本答案的其余部分并不是关于驱动程序 API 编程的教程（使用已经给出的参考资料和此处的API 参考手册），也不是关于 PTX 编程的教程。对于 PTX 编程，我建议您参阅PTX 文档。

首先，我们需要一个适当的 PTX 内核定义。（为此，我不会编写自己的内核 PTX 代码，而是使用vectorAddDrvCUDA 11.1 工具包中示例代码中的代码，通过将该 CUDA C++ 内核定义转换为等效的 PTX 内核定义nvcc -ptx vectorAdd_kernel.cu）：

矢量Add_kernel.ptx：

.version 7.1
.target sm_52
.address_size 64

        // .globl       VecAdd_kernel

.visible .entry VecAdd_kernel(
        .param .u64 VecAdd_kernel_param_0,
        .param .u64 VecAdd_kernel_param_1,
        .param .u64 VecAdd_kernel_param_2,
        .param .u32 VecAdd_kernel_param_3
)
{
        .reg .pred      %p<2>;
        .reg .f32       %f<4>;
        .reg .b32       %r<6>;
        .reg .b64       %rd<11>;


        ld.param.u64    %rd1, [VecAdd_kernel_param_0];
        ld.param.u64    %rd2, [VecAdd_kernel_param_1];
        ld.param.u64    %rd3, [VecAdd_kernel_param_2];
        ld.param.u32    %r2, [VecAdd_kernel_param_3];
        mov.u32         %r3, %ntid.x;
        mov.u32         %r4, %ctaid.x;
        mov.u32         %r5, %tid.x;
        mad.lo.s32      %r1, %r3, %r4, %r5;
        setp.ge.s32     %p1, %r1, %r2;
        @%p1 bra        $L__BB0_2;

        cvta.to.global.u64      %rd4, %rd1;
        mul.wide.s32    %rd5, %r1, 4;
        add.s64         %rd6, %rd4, %rd5;
        cvta.to.global.u64      %rd7, %rd2;
        add.s64         %rd8, %rd7, %rd5;
        ld.global.f32   %f1, [%rd8];
        ld.global.f32   %f2, [%rd6];
        add.f32         %f3, %f2, %f1;
        cvta.to.global.u64      %rd9, %rd3;
        add.s64         %rd10, %rd9, %rd5;
        st.global.f32   [%rd10], %f3;

$L__BB0_2:
        ret;

}

Run Code Online (Sandbox Code Playgroud)

我们还需要一个驱动程序 API C++ 源代码文件来完成所有主机端工作以加载该内核并启动它。vectorAddDrv我将再次使用示例项目（.cpp 文件）中的源代码，并进行修改以加载 PTX 而不是 fatbin：

矢量AddDrv.cpp：

// Vector addition: C = A + B.

// Includes
#include <stdio.h>
#include <string>
#include <iostream>
#include <cstring>
#include <fstream>
#include <streambuf>
#include <cuda.h>
#include <cmath>
#include <vector>

#define CHK(X) if ((err = X) != CUDA_SUCCESS) printf("CUDA error %d at %d\n", (int)err, __LINE__)

// Variables
CUdevice cuDevice;
CUcontext cuContext;
CUmodule cuModule;
CUfunction vecAdd_kernel;
CUresult err;
CUdeviceptr d_A;
CUdeviceptr d_B;
CUdeviceptr d_C;

// Host code
int main(int argc, char **argv)
{
    printf("Vector Addition (Driver API)\n");
    int N = 50000, devID = 0;
    size_t  size = N * sizeof(float);

    // Initialize
    CHK(cuInit(0));
    CHK(cuDeviceGet(&cuDevice, devID));
    // Create context
    CHK(cuCtxCreate(&cuContext, 0, cuDevice));
    // Load PTX file
    std::ifstream my_file("vectorAdd_kernel.ptx");
    std::string my_ptx((std::istreambuf_iterator<char>(my_file)), std::istreambuf_iterator<char>());
    // Create module from PTX
    CHK(cuModuleLoadData(&cuModule, my_ptx.c_str()));

    // Get function handle from module
    CHK(cuModuleGetFunction(&vecAdd_kernel, cuModule, "VecAdd_kernel"));

    // Allocate/initialize vectors in host memory
    std::vector<float> h_A(N, 1.0f);
    std::vector<float> h_B(N, 2.0f);
    std::vector<float> h_C(N);

    // Allocate vectors in device memory
    CHK(cuMemAlloc(&d_A, size));
    CHK(cuMemAlloc(&d_B, size));
    CHK(cuMemAlloc(&d_C, size));

    // Copy vectors from host memory to device memory
    CHK(cuMemcpyHtoD(d_A, h_A.data(), size));
    CHK(cuMemcpyHtoD(d_B, h_B.data(), size));

    // Grid/Block configuration
    int threadsPerBlock = 256;
    int blocksPerGrid   = (N + threadsPerBlock - 1) / threadsPerBlock;

    void *args[] = { &d_A, &d_B, &d_C, &N };

    // Launch the CUDA kernel
    CHK(cuLaunchKernel(vecAdd_kernel,  blocksPerGrid, 1, 1,
                               threadsPerBlock, 1, 1,
                               0,
                               NULL, args, NULL));

    // Copy result from device memory to host memory
    // h_C contains the result in host memory
    CHK(cuMemcpyDtoH(h_C.data(), d_C, size));

    // Verify result
    for (int i = 0; i < N; ++i)
    {
        float sum = h_A[i] + h_B[i];
        if (fabs(h_C[i] - sum) > 1e-7f)
        {
            printf("mismatch!");
            break;
        }
    }
    return 0;
}

Run Code Online (Sandbox Code Playgroud)

（请注意，我已经删除了各种项目，例如释放调用。这是为了演示整体方法；上面的代码只是一个演示器。）

在 Linux 上：

我们可以编译并运行代码如下：

$ g++ vectorAddDrv.cpp -o vectorAddDrv -I/usr/local/cuda/include -L/usr/local/cuda/lib64 -lcuda
$ ./vectorAddDrv
Vector Addition (Driver API)
$

Run Code Online (Sandbox Code Playgroud)

在 Windows/Visual Studio 上：在 Visual Studio 中创建一个新的 C++ 项目。将上面的.cpp 文件添加到项目中。确保该vectorAdd_kernel.ptx文件与构建的可执行文件位于同一目录中。您还需要修改项目定义以指向 CUDA 包含文件和 CUDA 库文件的位置。这是我在 VS2019 中所做的：

文件...新建...项目...控制台应用程序...创建
将给定 .cpp 文件的内容替换为上面的 .cpp 文件内容
将项目目标更改为 x64
在项目...属性中

将平台更改为 x64
在配置属性...C/C++...常规...附加包含目录中，添加CUDA工具包包含目录的路径，在我的机器上是C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.1\include
在配置属性...链接器...常规...附加库目录中，添加CUDA工具包库目录的路径，在我的机器上是C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.1\lib\x64
在配置属性...链接器...输入...附加依赖项中，添加文件cuda.lib（用于驱动程序 API 库）

保存项目属性，然后执行 Build....Rebuild
从控制台输出中，找到构建的可执行文件的位置。确保该vectorAdd_kernel.ptx文件位于该目录中，并从该目录运行可执行文件。（即打开命令提示符。更改到该目录。从命令提示符运行应用程序）

注意：如果您没有使用 CUDA 11.1 工具包或更高版本，或者如果您在计算能力 5.0 或更低的 GPU 上运行，则上述 PTX 代码将无法工作，因此此示例将无法逐字工作。然而整体方法是可行的，这个问题不是关于如何编写 PTX 代码。

编辑：回答评论中的问题：

如果您希望二进制文件不必在运行时构建任何内容怎么办？即组装 PTX 并将其与编译后的主机端代码粘贴到二进制文件中？

我不知道 NVIDIA 工具链提供了执行此操作的方法。从我的角度来看，创建这些统一的二进制文件几乎是运行时 API 的职责。

然而，从上面示例中的驱动程序 API 流程中可以看出，基本过程似乎很明显：无论我们从 a.cubin还是.ptx文件开始，无论哪种方式，文件都会加载到字符串中，并且字符串会被传递给cuModuleLoad()。.cubin因此，使用实用程序从二进制文件构建字符串，然后将其合并到构建过程中似乎并不困难。

我真的只是在这里进行黑客攻击，您应该自行承担使用它的风险，并且可能有很多我没有考虑到的因素。这部分我将在 Linux 上进行演示。以下是该实用程序的源代码和构建示例：

$ cat f2s.cpp
// Includes
#include <stdio.h>
#include <string>
#include <iostream>
#include <cstring>
#include <fstream>
#include <streambuf>

int main(int argc, char **argv)
{
    std::ifstream my_file("vectorAdd_kernel.cubin");
    std::string my_bin((std::istreambuf_iterator<char>(my_file)), std::istreambuf_iterator<char>());
    std::cout << "unsigned char my_bin[] = {";
    for (int i = 0; i < my_bin.length()-1; i++) std::cout << (int)(unsigned char)my_bin[i] << ",";
    std::cout << (int)(unsigned char)my_bin[my_bin.length()-1] << "};";
    return 0;
}
$ g++ f2s.cpp -o f2s
$

Run Code Online (Sandbox Code Playgroud)

下一步是创建一个.cubin供使用的文件。在上面的示例中，我通过nvcc -ptx vectorAdd_kernel.cu. 我们可以将其更改为，nvcc -cubin vectorAdd_kernel.cu或者您可以使用您喜欢的任何方法来生成文件.cubin。

创建 cubin 文件后，我们需要将其转换为可以纳入 C++ 代码构建过程的文件。这就是该实用程序的目的f2s。你可以像这样使用它：

./f2s > my_bin.h

Run Code Online (Sandbox Code Playgroud)

（可能最好允许该f2s实用程序接受输入文件名作为命令行参数。练习留给读者。这仅用于演示/娱乐。）

创建上述头文件后，我们需要修改我们的.cpp文件如下：

$ cat vectorAddDrv_bin.cpp
// Vector addition: C = A + B.

// Includes
#include <stdio.h>
#include <string>
#include <iostream>
#include <cstring>
#include <fstream>
#include <streambuf>
#include <cuda.h>
#include <cmath>
#include <vector>
#include <my_bin.h>
#define CHK(X) if ((err = X) != CUDA_SUCCESS) printf("CUDA error %d at %d\n", (int)err, __LINE__)

// Variables
CUdevice cuDevice;
CUcontext cuContext;
CUmodule cuModule;
CUfunction vecAdd_kernel;
CUresult err;
CUdeviceptr d_A;
CUdeviceptr d_B;
CUdeviceptr d_C;

// Host code
int main(int argc, char **argv)
{
    printf("Vector Addition (Driver API)\n");
    int N = 50000, devID = 0;
    size_t  size = N * sizeof(float);

    // Initialize
    CHK(cuInit(0));
    CHK(cuDeviceGet(&cuDevice, devID));
    // Create context
    CHK(cuCtxCreate(&cuContext, 0, cuDevice));
    // Create module from "binary string"
    CHK(cuModuleLoadData(&cuModule, my_bin));

    // Get function handle from module
    CHK(cuModuleGetFunction(&vecAdd_kernel, cuModule, "VecAdd_kernel"));

    // Allocate/initialize vectors in host memory
    std::vector<float> h_A(N, 1.0f);
    std::vector<float> h_B(N, 2.0f);
    std::vector<float> h_C(N);

    // Allocate vectors in device memory
    CHK(cuMemAlloc(&d_A, size));
    CHK(cuMemAlloc(&d_B, size));
    CHK(cuMemAlloc(&d_C, size));

    // Copy vectors from host memory to device memory
    CHK(cuMemcpyHtoD(d_A, h_A.data(), size));
    CHK(cuMemcpyHtoD(d_B, h_B.data(), size));

    // Grid/Block configuration
    int threadsPerBlock = 256;
    int blocksPerGrid   = (N + threadsPerBlock - 1) / threadsPerBlock;

    void *args[] = { &d_A, &d_B, &d_C, &N };

    // Launch the CUDA kernel
    CHK(cuLaunchKernel(vecAdd_kernel,  blocksPerGrid, 1, 1,
                               threadsPerBlock, 1, 1,
                               0,
                               NULL, args, NULL));

    // Copy result from device memory to host memory
    // h_C contains the result in host memory
    CHK(cuMemcpyDtoH(h_C.data(), d_C, size));

    // Verify result
    for (int i = 0; i < N; ++i)
    {
        float sum = h_A[i] + h_B[i];
        if (fabs(h_C[i] - sum) > 1e-7f)
        {
         printf("mismatch!");
            break;
        }
    }
    return 0;
}
$ g++ vectorAddDrv_bin.cpp -o vectorAddDrv_bin -I/usr/local/cuda/include -L/usr/local/cuda/lib64 -lcuda -I.
$ ./vectorAddDrv_bin
Vector Addition (Driver API)
$

Run Code Online (Sandbox Code Playgroud)

似乎有效。YMMV。为了进一步娱乐，这种方法似乎创建了一种混淆形式：

$ cuobjdump -sass vectorAdd_kernel.cubin

     code for sm_52
       Function : VecAdd_kernel
     .headerflags    @"EF_CUDA_SM52 EF_CUDA_PTX_SM(EF_CUDA_SM52)"
                                                                                 /* 0x001cfc00e22007f6 */
        /*0008*/                   MOV R1, c[0x0][0x20] ;                        /* 0x4c98078000870001 */
        /*0010*/                   S2R R0, SR_CTAID.X ;                          /* 0xf0c8000002570000 */
        /*0018*/                   S2R R2, SR_TID.X ;                            /* 0xf0c8000002170002 */
                                                                                 /* 0x001fd842fec20ff1 */
        /*0028*/                   XMAD.MRG R3, R0.reuse, c[0x0] [0x8].H1, RZ ;  /* 0x4f107f8000270003 */
        /*0030*/                   XMAD R2, R0.reuse, c[0x0] [0x8], R2 ;         /* 0x4e00010000270002 */
        /*0038*/                   XMAD.PSL.CBCC R0, R0.H1, R3.H1, R2 ;          /* 0x5b30011800370000 */
                                                                                 /* 0x001ff400fd4007ed */
        /*0048*/                   ISETP.GE.AND P0, PT, R0, c[0x0][0x158], PT ;  /* 0x4b6d038005670007 */
        /*0050*/                   NOP ;                                         /* 0x50b0000000070f00 */
        /*0058*/               @P0 EXIT ;                                        /* 0xe30000000000000f */
                                                                                 /* 0x081fd800fea207f1 */
        /*0068*/                   SHL R6, R0.reuse, 0x2 ;                       /* 0x3848000000270006 */
        /*0070*/                   SHR R0, R0, 0x1e ;                            /* 0x3829000001e70000 */
        /*0078*/                   IADD R4.CC, R6.reuse, c[0x0][0x140] ;         /* 0x4c10800005070604 */
                                                                                 /* 0x001fd800fe0207f2 */
        /*0088*/                   IADD.X R5, R0.reuse, c[0x0][0x144] ;          /* 0x4c10080005170005 */
        /*0090*/         {         IADD R2.CC, R6, c[0x0][0x148] ;               /* 0x4c10800005270602 */
        /*0098*/                   LDG.E R4, [R4]         }
                                                                                 /* 0xeed4200000070404 */
                                                                                 /* 0x001fd800f62007e2 */
        /*00a8*/                   IADD.X R3, R0, c[0x0][0x14c] ;                /* 0x4c10080005370003 */
        /*00b0*/                   LDG.E R2, [R2] ;                              /* 0xeed4200000070202 */
        /*00b8*/                   IADD R6.CC, R6, c[0x0][0x150] ;               /* 0x4c10800005470606 */
                                                                                 /* 0x001fc420fe4007f7 */
        /*00c8*/                   IADD.X R7, R0, c[0x0][0x154] ;                /* 0x4c10080005570007 */
        /*00d0*/                   FADD R0, R2, R4 ;                             /* 0x5c58000000470200 */
        /*00d8*/                   STG.E [R6], R0 ;                              /* 0xeedc200000070600 */
                                                                                 /* 0x001ffc00ffe007ea */
        /*00e8*/                   NOP ;                                         /* 0x50b0000000070f00 */
        /*00f0*/                   EXIT ;                                        /* 0xe30000000007000f */
        /*00f8*/                   BRA 0xf8 ;                                    /* 0xe2400fffff87000f */
       ..........


$ cuobjdump -sass vectorAddDrv_bin
cuobjdump info    : File 'vectorAddDrv_bin' does not contain device code
$

Run Code Online (Sandbox Code Playgroud)

哈哈

归档时间：	4 年，4 月前
查看次数：	1297 次
最近记录：	4 年，4 月前