如何编译PTX代码

Question

我需要修改PTX代码并直接编译它。原因是我希望有一些特定的指令在彼此之后，而且很难编写一个cuda代码，从而导致我的目标PTX代码，所以我需要直接修改ptx代码。问题是我可以将其编译到(fatbin和cubin)，但我不知道如何将这些文件(.fatbin和.cubin)编译为"X.o“文件。

Answer 1

也许有一种方法可以通过有序的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;
}

为此，我不再使用库达错误检查和其他细节，而倾向于简洁。

通常，我们可以按以下方式编译上述代码：

nvcc -arch=sm_20 -o t266 t266.cu 

(假设源文件名为t266.cu)

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

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

这将生成可执行文件，但将保留所有中间文件，包括t266.ptx (其中包含mykernel的ptx代码)。

如果我们现在只运行可执行文件，就会得到如下的输出：

$ ./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

(这里使用linux重定向stderr )。然后，我们希望编辑该dryrun.out文件，以便：

1. 在创建ptx文件之后，我们保留所有的命令，直到文件的末尾。创建ptx文件的行显然是指定-o "t266.ptx"的行。
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
$

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

Answer 2

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

Answer 3

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

您可以使用它将PTX包含到您的构建中，尽管这个方法有点复杂。例如，素丽卡将其.cu文件编译为针对不同体系结构的PTX文件，然后将它们转换为包含PTX代码的.h文件作为'C‘数组，然后在构建过程中将其从其中一个文件中包含进去。