10

我需要修改PTX代码并直接编译。原因是我想要一个接一个的特定指令，并且很难编写导致我的目标 PTX 代码的 cuda 代码，所以我需要直接修改 ptx 代码。问题是我可以将其编译为（fatbin 和 cubin），但我不知道如何将这些（.fatbin 和 .cubin）编译为“Xo”文件。

4

6 回答 6

10

可能有一种方法可以通过有序的nvcc命令序列来做到这一点，但我不知道也没有发现它。

然而，一种可能的方法（尽管很麻烦）是中断并重新启动 cuda 编译序列，并在中间（重新启动之前）编辑 ptx 文件。这是基于nvcc 手册中提供的信息，我不认为这是标准方法，因此您的里程可能会有所不同。可能有许多我没有考虑过的情况，这不起作用或不可行。

为了解释这一点，我将提供一个示例代码：

#include <stdio.h>

__global__ void mykernel(int *data){

  (*data)++;
}

int main(){

  int *d_data, h_data = 0;
  cudaMalloc((void **)&d_data, sizeof(int));
  cudaMemcpy(d_data, &h_data, sizeof(int), cudaMemcpyHostToDevice);
  mykernel<<<1,1>>>(d_data);
  cudaMemcpy(&h_data, d_data, sizeof(int), cudaMemcpyDeviceToHost);
  printf("data = %d\n", h_data);
  return 0;
}

为此，为了简洁，我放弃了cuda 错误检查和其他细节。

通常我们可以将上面的代码编译如下：

nvcc -arch=sm_20 -o t266 t266.cu

（假设源文件名为 t266.cu）

相反，根据参考手册，我们将编译如下：

nvcc -arch=sm_20 -keep -o t266 t266.cu

这将构建可执行文件，但将保留所有中间文件，包括t266.ptx（包含 ptx 代码mykernel）

如果我们此时简单地运行可执行文件，我们将得到如下输出：

$ ./t266
data = 1
$

下一步将是编辑 ptx 文件以进行我们想要的任何更改。在这种情况下，我们将让内核将 2 添加到data变量而不是添加 1。相关行是：

    add.s32         %r2, %r1, 2;
                              ^
                              |
                          change the 1 to a 2 here

现在是混乱的部分。下一步是捕获所有中间编译命令，因此我们可以重新运行其中一些命令：

nvcc -dryrun -arch=sm_20 -o t266 t266.cu --keep 2>dryrun.out

（使用stderr这里的 linux 重定向）。然后我们要编辑该dryrun.out文件，以便：

  1. 我们保留创建 ptx 文件后的所有命令，直到文件末尾。创建 ptx 文件的行将明显指定为-o "t266.ptx"
  2. 我们去掉了每一行开头的前导#$，所以实际上我们正在创建一个脚本。

当我执行上述 2 个步骤时，我最终会得到如下脚本：

ptxas  -arch=sm_20 -m64  "t266.ptx"  -o "t266.sm_20.cubin"
fatbinary --create="t266.fatbin" -64 --key="xxxxxxxxxx" --ident="t266.cu" "--image=profile=sm_20,file=t266.sm_20.cubin" "--image=profile=compute_20,file=t266.ptx" --embedded-fatbin="t266.fatbin.c" --cuda
gcc -D__CUDA_ARCH__=200 -E -x c++   -DCUDA_DOUBLE_MATH_FUNCTIONS   -D__CUDA_PREC_DIV -D__CUDA_PREC_SQRT "-I/usr/local/cuda/bin/..//include"   -m64 -o "t266.cu.cpp.ii" "t266.cudafe1.cpp"
gcc -c -x c++ "-I/usr/local/cuda/bin/..//include"   -fpreprocessed -m64 -o "t266.o" "t266.cu.cpp.ii"
nvlink --arch=sm_20 --register-link-binaries="t266_dlink.reg.c" -m64   "-L/usr/local/cuda/bin/..//lib64" "t266.o"  -o "t266_dlink.sm_20.cubin"
fatbinary --create="t266_dlink.fatbin" -64 --key="t266_dlink" --ident="t266.cu " -link "--image=profile=sm_20,file=t266_dlink.sm_20.cubin" --embedded-fatbin="t266_dlink.fatbin.c"
gcc -c -x c++ -DFATBINFILE="\"t266_dlink.fatbin.c\"" -DREGISTERLINKBINARYFILE="\"t266_dlink.reg.c\"" -I. "-I/usr/local/cuda/bin/..//include"   -m64 -o "t266_dlink.o" "/usr/local/cuda/bin/crt/link.stub"
g++ -m64 -o "t266" -Wl,--start-group "t266_dlink.o" "t266.o"   "-L/usr/local/cuda/bin/..//lib64" -lcudart_static  -lrt -lpthread -ldl  -Wl,--end-group

最后，执行上面的脚本。（在 linux 中，您可以使用或类似工具使该脚本文件可执行chmod +x dryrun.out。）如果您在编辑文件时没有犯任何错误.ptx，则命令应该全部成功完成，并创建一个新的t266可执行文件。

当我们运行该文件时，我们观察到：

$ ./t266
data = 2
$

表明我们的更改是成功的。

于 2013-11-16T04:07:09.707 回答
3

通常，在处理 cubin 或 ptx-files 时，使用 CUDA Driver API 而不是 Runtime API；这样做，您可以在运行时手动加载 ptx 或 cubin 文件cuModuleLoadDataEx。如果您想坚持使用运行时 API，您需要手动模仿 NVCC 所做的事情，但这没有（完全）记录在案。我只找到了这个关于如何做到这 一点的Nvidia 论坛条目。

于 2013-11-16T00:04:47.667 回答
0

我来晚了，但GPU Lynx正是这样做的：获取一个 CUDA 胖二进制文件，解析 PTX，并在将结果发送给驱动程序以在 GPU 上执行之前对其进行修改。您也可以选择打印出修改后的 PTX。

于 2015-04-09T18:30:40.237 回答
0

您可以在运行时使用 CUDA 中的 cuModuleLoad* 函数加载 cubin 或 fatbin：这是 API

您可以使用它将 PTX 包含到您的构建中，尽管该方法有些复杂。例如，suricata将其 .cu 文件编译为不同架构的 PTX 文件，然后将它们转换为 .h 文件，其中包含作为“C”数组的 PTX 代码，然后在构建期间仅从其中一个文件中包含它。

于 2013-11-16T00:04:38.300 回答
0

这个 nvcc 命令序列似乎可以解决问题。请参阅此处了解更多详情。

创建要修改的 ptx 文件

nvcc file1.cu file2.cu file3.cu -rdc=true --ptx

将 ptx 文件链接到目标文件

nvcc file1.ptx file2.ptx file3.ptx -dlink

我在 Windows 上做了这个，所以它弹出了a_dlink.obj。正如文档指出的那样，此时主机代码已被丢弃。跑

nvcc file1.cu file2.cu file3.cu -rdc=true --compile

创建目标文件。它们将.obj适用于 Windows 或.oLinux。然后创建一个库输出文件

nvcc file1.obj file2.obj file3.obj a_dlink.obj --lib -o myprogram.lib

然后运行

nvcc myprogram.lib

a.exe这将在 Windows 或a.outLinux上弹出一个可执行文件。此过程也适用于cubin和fatbin文件。只需将这些名称替换为ptx.

于 2017-09-29T06:55:11.930 回答
0

您可以使用 NVTRC - 这很简单！

扩展@ArtemB的答案：

nVIDIA 提供实时编译 (RTC) 库。有一个如何将其用作 CUDA 示例的示例；你可以在这里访问它。

该示例实际上从 CUDA 代码开始，但中间步骤是将 PTX 代码创建为纯 C 字符串 (`char *)。从那里开始，这就是你所做的，基本上：

char* ptx;
size_t ptxSize;

// ... populate ptx and ptxSize somehow ...

CUcontext context;
CUdevice cuDevice;

// These next few lines simply initialize your work with the CUDA driver,
// they're not specific to PTX compilation
cuInit(0);
cuDeviceGet(&cuDevice, 0); // or some other device on your system
cuCtxCreate(&context, 0, cuDevice);

// The magic happens here:
CUmodule module;
cuModuleLoadDataEx(&module, ptx, 0, 0, 0));

// And here is how you use your compiled PTX
CUfunction kernel_addr;
cuModuleGetFunction(&kernel_addr, module, "my_kernel_name");
cuLaunchKernel(kernel_addr, 
   // launch parameters go here
   // kernel arguments go here
);

笔记：

  • 我已删除所有错误检查，以免弄乱示例代码 - 但请检查代码中的错误。
  • 您需要将程序与 NVRTC 库链接 - 它与主 CUDA 和 CUDA 驱动程序库分开。在 linux 上，它被称为libnvrtc.so.
于 2020-01-22T22:49:47.180 回答