如何在给定 PTX 文件中为 运行 内核创建可执行文件？
How can I create an executable to run a kernel in a given PTX file?
据我所知，您需要一个主机代码（针对 CPU）和一个设备代码（针对 GPU），没有它们您就无法 运行 GPU 上的东西。
我正在学习 PTX ISA，但我不知道如何在 Windows 上执行它。我需要一个 .cu 文件来 运行 它还是有另一种方法来 运行 它？
TL;DR:
How can I assemble .ptx file and host code file and make a executable file?
您使用CUDA driver API. Relevant sample codes are vectorAddDrv (or perhaps any other driver API sample code) as well as ptxjit.
Do I need a .cu file to run it or is there another way to run it?
如果您从 PTX 格式的设备代码开始，则不需要 .cu 文件（也不需要 nvcc）来使用驱动程序 API 方法。
详情：
这个答案的其余部分并不是驱动程序 API 编程的教程（使用已经给出的参考文献和 API 参考手册 here), nor is it intended to be a tutorial on PTX programming. For PTX programming I refer you to the PTX documentation.
首先，我们需要一个适当的 PTX 内核定义。 （为此，我不会编写自己的内核 PTX 代码，而是使用 CUDA 11.1 工具包中 vectorAddDrv 示例代码中的代码，通过 nvcc -ptx vectorAdd_kernel.cu):
vectorAdd_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;
}
我们还需要一个驱动程序 API C++ 源代码文件来完成所有主机端工作以加载并启动该内核。我将再次使用 vectorAddDrv 示例项目（.cpp 文件）中的源代码，并进行修改以加载 PTX 而不是 fatbin：
vectorAddDrv.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;
}
（请注意，我已经删除了各种项目，例如释放调用。这是为了演示整体方法；上面的代码只是一个演示。）
在 Linux:
我们可以编译运行代码如下：
$ g++ vectorAddDrv.cpp -o vectorAddDrv -I/usr/local/cuda/include -L/usr/local/cuda/lib64 -lcuda
$ ./vectorAddDrv
Vector Addition (Driver API)
$
在 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
- 在配置属性...Linker...General...Additional Library Directories中，添加CUDA工具包库目录的路径，在我的机器上是
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.1\lib\x64
- 在配置属性...链接器...输入...附加依赖项中，添加
cuda.lib 文件（用于驱动程序 API 库）
- 保存项目属性，然后执行构建....重建
- 从控制台输出中，找到构建的可执行文件的位置。确保
vectorAdd_kernel.ptx 文件在该目录中，并且 运行 该目录中的可执行文件。 （即打开命令提示符。更改到该目录。运行 来自命令提示符的应用程序）
注意：如果您没有使用 CUDA 11.1 工具包或更新版本，或者如果您 运行正在使用计算能力为 5.0 或更低的 GPU，则上述 PTX 代码将不起作用，因此此示例将不会逐字工作。不过整体的方法是可以的，这个问题不是PTX代码怎么写。
编辑：回答评论中的问题：
What if you wanted the binary to not have to build anything at runtime? i.e. assemble the PTX and stick it in a binary with the compiled host-side code?
我不知道 NVIDIA 工具链提供了执行此操作的方法。从我的角度来看，创建这些统一的二进制文件几乎是 domain of the runtime API。
然而，从上面示例中的驱动程序 API 流程可以看出，基本过程似乎很明显：我们是从 .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
$
这里的下一步是创建一个 .cubin 文件以供使用。在上面的示例中，我通过 nvcc -ptx vectorAdd_kernel.cu 创建了 ptx 文件。我们可以将其更改为 nvcc -cubin vectorAdd_kernel.cu 或者您可以使用任何您喜欢的方法来生成 .cubin 文件。
创建 cubin 文件后，我们需要将其转换为可以纳入我们的 C++ 代码构建过程的内容。这就是 f2s 实用程序的目的。你会像这样使用它：
./f2s > my_bin.h
（可能允许 f2s 实用程序接受输入文件名作为命令行参数会很好。练习留给 reader。这仅适用于 demonstration/amusement .)
创建完上面的头文件后，我们需要修改我们的.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)
$
这似乎有效。 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
$
英雄联盟
据我所知，您需要一个主机代码（针对 CPU）和一个设备代码（针对 GPU），没有它们您就无法 运行 GPU 上的东西。
我正在学习 PTX ISA，但我不知道如何在 Windows 上执行它。我需要一个 .cu 文件来 运行 它还是有另一种方法来 运行 它？
TL;DR:
How can I assemble .ptx file and host code file and make a executable file?
您使用CUDA driver API. Relevant sample codes are vectorAddDrv (or perhaps any other driver API sample code) as well as ptxjit.
Do I need a .cu file to run it or is there another way to run it?
如果您从 PTX 格式的设备代码开始，则不需要 .cu 文件（也不需要 nvcc）来使用驱动程序 API 方法。
详情：
这个答案的其余部分并不是驱动程序 API 编程的教程（使用已经给出的参考文献和 API 参考手册 here), nor is it intended to be a tutorial on PTX programming. For PTX programming I refer you to the PTX documentation.
首先，我们需要一个适当的 PTX 内核定义。 （为此，我不会编写自己的内核 PTX 代码，而是使用 CUDA 11.1 工具包中 vectorAddDrv 示例代码中的代码，通过 nvcc -ptx vectorAdd_kernel.cu):
vectorAdd_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;
}
我们还需要一个驱动程序 API C++ 源代码文件来完成所有主机端工作以加载并启动该内核。我将再次使用 vectorAddDrv 示例项目（.cpp 文件）中的源代码，并进行修改以加载 PTX 而不是 fatbin：
vectorAddDrv.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;
}
（请注意，我已经删除了各种项目，例如释放调用。这是为了演示整体方法；上面的代码只是一个演示。）
在 Linux:
我们可以编译运行代码如下：
$ g++ vectorAddDrv.cpp -o vectorAddDrv -I/usr/local/cuda/include -L/usr/local/cuda/lib64 -lcuda
$ ./vectorAddDrv
Vector Addition (Driver API)
$
在 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 - 在配置属性...Linker...General...Additional Library Directories中，添加CUDA工具包库目录的路径，在我的机器上是
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.1\lib\x64 - 在配置属性...链接器...输入...附加依赖项中，添加
cuda.lib文件（用于驱动程序 API 库）
- 保存项目属性，然后执行构建....重建
- 从控制台输出中，找到构建的可执行文件的位置。确保
vectorAdd_kernel.ptx文件在该目录中，并且 运行 该目录中的可执行文件。 （即打开命令提示符。更改到该目录。运行 来自命令提示符的应用程序）
注意：如果您没有使用 CUDA 11.1 工具包或更新版本，或者如果您 运行正在使用计算能力为 5.0 或更低的 GPU，则上述 PTX 代码将不起作用，因此此示例将不会逐字工作。不过整体的方法是可以的，这个问题不是PTX代码怎么写。
编辑：回答评论中的问题：
What if you wanted the binary to not have to build anything at runtime? i.e. assemble the PTX and stick it in a binary with the compiled host-side code?
我不知道 NVIDIA 工具链提供了执行此操作的方法。从我的角度来看，创建这些统一的二进制文件几乎是 domain of the runtime API。
然而，从上面示例中的驱动程序 API 流程可以看出，基本过程似乎很明显：我们是从 .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
$
这里的下一步是创建一个 .cubin 文件以供使用。在上面的示例中，我通过 nvcc -ptx vectorAdd_kernel.cu 创建了 ptx 文件。我们可以将其更改为 nvcc -cubin vectorAdd_kernel.cu 或者您可以使用任何您喜欢的方法来生成 .cubin 文件。
创建 cubin 文件后，我们需要将其转换为可以纳入我们的 C++ 代码构建过程的内容。这就是 f2s 实用程序的目的。你会像这样使用它：
./f2s > my_bin.h
（可能允许 f2s 实用程序接受输入文件名作为命令行参数会很好。练习留给 reader。这仅适用于 demonstration/amusement .)
创建完上面的头文件后，我们需要修改我们的.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)
$
这似乎有效。 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
$
英雄联盟