如何编译PTX代码

Question

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

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;
}

为此,我不再需要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. 我们删除#$每行开头的前导,所以实际上我们正在创建一个脚本.

当我执行上述两个步骤时,我最终得到一个这样的脚本:

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

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

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