有没有办法访问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

Answer 2

接受的解决方案对我不起作用。我可以执行以下操作：

(cuda-gdb) disass $pc,+16
Dump of assembler code from 0x7fffd5043d40 to 0x7fffd5043d50:
=> 0x00007fffd5043d40 <_Z9acos_main10acosParams+1856>:  MOV R0, c[0x0][0xc]
End of assembler dump.
(cuda-gdb) p *$_cuda_const_bank(0x0,0xc)
$1 = 8

这来自 cuda-gdb 参考。