微信公众号搜"智元新知"关注
微信扫一扫可直接关注哦！

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

分类：编程问答

如何解决如何创建可执行文件以在给定的 PTX 文件中运行内核？

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

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

解决方法

TL;博士：

如何组合 .ptx 文件和主机代码文件并制作可执行文件？

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

我需要一个 .cu 文件来运行它还是有其他方法来运行它？

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

详情：

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

首先，我们需要一个合适的 PTX 内核定义。 （为此，我将使用 CUDA 11.1 工具包中 vectorAddDrv 示例代码中的代码，而不是编写自己的内核 PTX 代码，通过 {{1} 将该 CUDA C++ 内核定义转换为等效的 PTX 内核定义}}):

vectorAdd_kernel.ptx:

nvcc -ptx vectorAdd_kernel.cu

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

vectorAddDrv.cpp:

vectorAddDrv

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

在 Linux 上：

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

// 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,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,threadsPerBlock,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;
}

在 Windows/Visual Studio 上：在 Visual Studio 中创建一个新的 C++ 项目。将上述 .cpp 文件添加到项目中。确保 $ g++ vectorAddDrv.cpp -o vectorAddDrv -I/usr/local/cuda/include -L/usr/local/cuda/lib64 -lcuda $ ./vectorAddDrv Vector Addition (Driver API) $ 文件与构建的可执行文件位于同一目录中。您还需要修改项目定义以指向 CUDA 包含文件和 CUDA 库文件的位置。这是我在 VS2019 中所做的：

  1. 文件...新建...项目...控制台应用...创建
  2. 用上面的 .cpp 文件内容替换给定 .cpp 文件的内容
  3. 将项目目标更改为 x64
  4. 在项目中...属性
  • 将平台更改为 x64
  • 在配置属性...C/C++...General...Additional Include Directories 中，添加到 CUDA 工具包包含目录的路径，在我的机器上它是 vectorAdd_kernel.ptx
  • 在配置属性...链接器...常规...附加库目录中，添加到 CUDA 工具包库目录的路径，在我的机器上它是 C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.1\include
  • 在配置属性...Linker...Input...Additional Dependencies 中，添加 C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.1\lib\x64 文件（用于驱动程序 API 库）
  1. 保存项目属性，然后执行 Build....Rebuild
  2. 从控制台输出中，找到构建的可执行文件的位置。确保 cuda.lib 文件在该目录中，并从该目录运行可执行文件。 （即打开命令提示符。更改到该目录。从命令提示符运行应用程序）

注意：如果您没有使用 CUDA 11.1 或更高版本的工具包，或者您在计算能力为 5.0 或更低的 GPU 上运行，则上述 PTX 代码将不起作用，因此该示例将无法逐字运行。不过整体方法还是行得通的，这个问题不是怎么写PTX代码。

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

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

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

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

我真的只是在这里闲聊，你应该自担风险使用它，而且可能有许多我没有考虑的因素。这部分我只是要在 linux 上进行演示。这是该实用程序的源代码和构建示例：

.cubin

此处的下一步是创建一个 $ 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 $ 文件以供使用。在上面的例子中，我通过 .cubin 创建了 ptx 文件。我们可以将其更改为 nvcc -ptx vectorAdd_kernel.cu，或者您可以使用任何您喜欢的方法来生成 nvcc -cubin vectorAdd_kernel.cu 文件。

创建 cubin 文件后，我们需要将其转换为可以吸收到我们的 C++ 代码构建过程中的内容。这就是 .cubin 实用程序的目的。你会像这样使用它：

f2s

（可能允许 ./f2s > my_bin.h 实用程序接受输入文件名作为命令行参数会很好。练习留给读者。这只是为了演示/娱乐。）

上述头文件创建完成后，我们需要修改我们的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",cuDevice));
    // Create module from "binary string"
    CHK(cuModuleLoadData(&cuModule,my_bin));

    // Get function handle from module
    CHK(cuModuleGetFunction(&vecAdd_kernel,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)
$

大声笑

版权声明：本文内容由互联网用户自发贡献，该文观点与技术仅代表作者本人。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如发现本站有涉嫌侵权/违法违规的内容， 请发送邮件至 dio@foxmail.com 举报，一经查实，本站将立刻删除。

苹果市值2025年有望达4万亿美元
pythonJavaScriptjavaHTMLPHPreactjsC#AndroidCSSNode.jssqlrpython-3.xMysqLjQueryc++pandasFlutterangularIOSdjangolinuxswifttypescript路由器JSON路由器设置无线路由器h3c华三华三路由器设置华三路由器电脑软件教程arraysdocker软件图文教程Cvue.jslaravelspring-boot