有没有办法在 CUDA 中访问常量内存库的值

Question

我一直在尝试调试使用内联 PTX 程序集的 cuda 程序。具体来说，我在指令级别进行调试，并试图确定指令参数的值。有时，反汇编包含对常量内存的引用。我正在尝试让 gdb 打印此常量内存的值，但没有找到任何说明如何执行此操作的文档。 例如，反汇编包括 IADD R0, R0, c[0x0] [0x148]

我想确定如何让 gdb 打印 c[0x0] [0x148] 的值。我试过使用 print * (@constant) ... 但这似乎不起作用（我在这里传递了 0x148 并且它什么也没打印出来）。这可以在 cuda-gdb 中做到吗？

我试图通过在编译期间传递编译器选项 --disable-optimizer-constants 来避免这种情况，但这不起作用。

Answer 1

这样做的方法是 print *(void * @parameter *) addr

其中 addr 是常量 bank 0 中应打印的地址。

示例

假设我们在一个名为 foo.cu 的文件中有一个简单的内核：

#include <cuda.h>
#include <stdio.h>
#include <cuda_runtime.h>

__global__ void myKernel(int a, int b, int *d)
{
    *d = a + b;
}

int main(int argc, char *argv[]) {
   if (argc < 3) {
       printf("Requires inputs a and b to be specified\n");
       return 0;
   }

   int * dev_d;
   int d;
   cudaMalloc(&dev_d, sizeof(*dev_d));
   myKernel<<<1, 1>>>(atoi(argv[1]), atoi(argv[2]), dev_d);
   cudaMemcpy(&d, dev_d, sizeof(d), cudaMemcpyDeviceToHost);
   cudaFree(dev_d);
   printf("D is: %d\n", d);
   return 0;
}

编译通过

$ nvcc foo.cu -o foo.out

接下来，假设我们有兴趣反汇编这个程序，所以我们在程序的命令行中执行cuda-gdb：

$ cuda-gdb --args ./foo.out 10 15

在cuda-gdb 中，我们通过键入进入内核

(cuda-gdb) set cuda break_on_launch application
(cuda-gdb) start
Temporary breakpoint 1, 0x000055555555b12a in main ()

(cuda-gdb) cont

在内核内部，我们查看我们有兴趣调试的反汇编：

(cuda-gdb) x/15i $pc
=> 0x555555b790a8 <_Z8myKerneliiPi+8>:  MOV R1, c[0x0][0x20]
   0x555555b790b0 <_Z8myKerneliiPi+16>: MOV R0, c[0x0][0x144]
   0x555555b790b8 <_Z8myKerneliiPi+24>: MOV R2, c[0x0][0x148]
   0x555555b790c0 <_Z8myKerneliiPi+32>:
   0x555555b790c8 <_Z8myKerneliiPi+40>: MOV R3, c[0x0][0x14c]
   0x555555b790d0 <_Z8myKerneliiPi+48>: IADD R0, R0, c[0x0][0x140]
   0x555555b790d8 <_Z8myKerneliiPi+56>: STG.E [R2], R0
   0x555555b790e0 <_Z8myKerneliiPi+64>:
   0x555555b790e8 <_Z8myKerneliiPi+72>: NOP
   0x555555b790f0 <_Z8myKerneliiPi+80>: NOP
   0x555555b790f8 <_Z8myKerneliiPi+88>: NOP
   0x555555b79100 <_Z8myKerneliiPi+96>:
   0x555555b79108 <_Z8myKerneliiPi+104>:        EXIT
   0x555555b79110 <_Z8myKerneliiPi+112>:        BRA 0x70
   0x555555b79118 <_Z8myKerneliiPi+120>:        NOP

传递给IADD 指令的第二个参数位于其中一个常量内存库中。让我们找出它的实际价值。我们提前转到IADD 指令：

(cuda-gdb) stepi 4
0x0000555555b790d0 in myKernel(int, int, int*)<<<(1,1,1),(1,1,1)>>> ()
(cuda-gdb) x/i $pc
=> 0x555555b790d0 <_Z8myKerneliiPi+48>: IADD R0, R0, c[0x0][0x140]

我们现在可以得到c[0x0][0x140]的内容如下：

(cuda-gdb) print (int) *(void * @parameter *) 0x140
$1 = 10

在这里，我们知道参数应该有 32 位，因此我们将其转换为（32 位）int。如果我们没有这样做，我们会得到太多位，例如：

(cuda-gdb) print *(void * @parameter *) 0x140
$2 = 0xf0000000a

注意十六进制格式可以通过在print命令后加/x来保留：

(cuda-gdb) print/x (int) *(void * @parameter *)0x140
$3 = 0xa