Cuda gdb print constant

Question

I am in cuda-gdb, I can use ((@global float *)array)[0]

but how to use constant memory in gdb ?

I try ((@parameter float *)const_array)

I declared const_array like this : __constant__ float const_array[1 << 14]

I tried with 1 << 5, and it's the same problem.

Answer 1

Try putting you __constant__ into .cuh, then use as a classic C global variable.

Answer 2

I don't seem to have any trouble with it. In order to print device memory, you must be stopped at a breakpoint in device code.

Example:

$ cat t1973.cu
const int cs = 1 << 14;
__constant__ int cdata[cs];
__global__ void k(int *gdata){

  gdata[0] = cdata[0];
}

int main(){

  int *hdata = new int[cs];
  for (int i = 0; i < cs; i++) hdata[i] = i+1;
  cudaMemcpyToSymbol(cdata, hdata, cs*sizeof(cdata[0]));
  int *gdata;
  cudaMalloc(&gdata, sizeof(gdata[0]));
  cudaMemset(gdata, 0, sizeof(gdata[0]));
  k<<<1,1>>>(gdata);
  cudaDeviceSynchronize();
}
$ nvcc -o t1973 t1973.cu -g -G -arch=sm_70
$ cuda-gdb ./t1973
sh: python3: command not found
Unable to determine python3 interpreter version. Python integration disabled.
NVIDIA (R) CUDA Debugger
11.4 release
Portions Copyright (C) 2007-2021 NVIDIA Corporation
GNU gdb (GDB) 10.1
Copyright (C) 2020 Free Software Foundation, Inc.
License GPLv3+: GNU GPL version 3 or later <http://gnu.org/licenses/gpl.html>
This is free software: you are free to change and redistribute it.
There is NO WARRANTY, to the extent permitted by law.
Type "show copying" and "show warranty" for details.
This GDB was configured as "x86_64-pc-linux-gnu".
Type "show configuration" for configuration details.
For bug reporting instructions, please see:
<https://www.gnu.org/software/gdb/bugs/>.
Find the GDB manual and other documentation resources online at:
    <http://www.gnu.org/software/gdb/documentation/>.

For help, type "help".
Type "apropos word" to search for commands related to "word"...
Reading symbols from ./t1973...
(cuda-gdb) b 5
Breakpoint 1 at 0x403b0c: file t1973.cu, line 6.
(cuda-gdb) run
Starting program: /home/user2/misc/t1973
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib64/libthread_db.so.1".
[Detaching after fork from child process 22872]
[New Thread 0x7fffef475700 (LWP 22879)]
[New Thread 0x7fffeec74700 (LWP 22880)]
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]

Thread 1 "t1973" hit Breakpoint 1, k<<<(1,1,1),(1,1,1)>>> (
    gdata=0x7fffcdc00000) at t1973.cu:5
5         gdata[0] = cdata[0];
(cuda-gdb) print gdata[0]
$1 = 0
(cuda-gdb) print cdata[0]
$2 = 1
(cuda-gdb) s
6       }
(cuda-gdb) print gdata[0]
$3 = 1
(cuda-gdb) print cdata[0]
$4 = 1
(cuda-gdb) print cdata[1]
$5 = 2
(cuda-gdb)