cuda - 如何编译PTX代码

Question

我需要修改PTX代码并直接进行编译。原因是我想彼此之间有一些特定的说明，并且很难编写生成目标PTX代码的cuda代码，因此我需要直接修改ptx代码。
问题是我可以将其编译为(fatbin和cubin)，但是我不知道如何将它们(.fatbin和.cubin)编译为“X.o”文件。

Answer 1

也许可以通过按顺序排列的nvcc命令来执行此操作，但是我不知道并且也没有发现它。

但是，尽管很麻烦，但是一种可能的方法是中断并重新启动cuda编译序列，并在过渡期间(重新启动之前)编辑ptx文件。这基于nvcc manual中提供的信息，我不认为这是标准方法，因此您的里程可能会有所不同。在这种情况行不通或不可行的情况下，我可能没有考虑过多种情况。

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

#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 error checking和其他功能。

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

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文件，以便:

我们将在创建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
$

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