Cuda-gdb prints struct member values incorrectly

Developer Tools CUDA Developer Tools CUDA-GDB
1 
#include <stdio.h>

#define CHECK_CUDA(call)                                                       \
  {                                                                            \
    cudaError_t err = call;                                                    \
    if (cudaSuccess != err) {                                                  \
      fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n", __FILE__,  \
              __LINE__, cudaGetErrorString(err));                              \
      exit(EXIT_FAILURE);                                                      \
    }                                                                          \
  }


struct d{
    int s;
    int e;
    int id;
};
__global__ void test()
{
    d da{0, 256, 192};
    printf("%d,%d,%d\n", da.s, da.e, da.id);
}


int main()
{
    test<<<1,256>>>();
    CHECK_CUDA(cudaDeviceSynchronize());
}

In this example, if I set a breakpoint to the printf statement in the kernel and print the struct instance da

(cuda-gdb) p da
$1 = {s = 0, e = 256, id = 192}

I get the correct result

But if I print only one element I would get incorrect result

(cuda-gdb) p da.id
$2 = 2013397007
(cuda-gdb) p da.e
$3 = -469760013

Why can’t I query each member of the struct like this and get the correct value?

cuda-gdb version :

cuda-gdb --version
NVIDIA (R) CUDA Debugger
11.5 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.
2

Hi @BAdhi
Thank you for your report we are looking at the issue. Could you also share a few more details:

  • nvcc compilation command used to build the repro case
  • nvcc --version output
3

Hi @AKravets ,

compile command:

 nvcc ./test.cu -g -G -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80 -t 2

nvcc version:

nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2021 NVIDIA Corporation
Built on Mon_Sep_13_19:13:29_PDT_2021
Cuda compilation tools, release 11.5, V11.5.50
Build cuda_11.5.r11.5/compiler.30411180_0
4

Thanks for the reproducer! I can confirm that I am seeing the same behavior internally and we are working on a fix to address this. It will be addressed in a future release. Thanks much for the report.

6

The cudbgprocess is an RPC interface belonging to the cuda driver. We communicate with this process via pipes. You can see this code from the cuda-gdb side primarily in gdb/cuda/libcudbgipc.c and gdb/cuda/libcudbg.c.

1 Like
7

Is there a work around to get the correct value? I need to set conditional breakpoints based on the member value and It doesn’t seem to work due to this error

8

I can confirm that this issue has been resolved. The fix will be included with the cuda-gdb found with the next major release of the CUDA Toolkit.

(cuda-gdb) p da
$11 = {s = 0, e = 256, id = 192}
(cuda-gdb) p da.s
$12 = 0
(cuda-gdb) p da.e
$13 = 256
(cuda-gdb) p da.id
$14 = 192
1 Like
9

I hit this issue, interestingly exactly one year after this post.
My CUDA toolkit is exactly same version as the post, and are there any workaround to debug the values of struct without updating CUDA toolkit?