cuda-gdb loses track of variables when stepping into functions..

I get “warning: Variable is not live at this point. Returning garbage value.” in cuda-gdb while examing a variable pointer that was passed as parameter.

Basically I do this:

[font=“Courier New”]float a;


Then when stepping into [font=“Courier New”]foo()[/font], “[font=“Courier New”]print a[/font]” and “[font=“Courier New”]print *a[/font]” give a warning.

The local variable can also not be printed after returning from the function.

The following program works correctly and prints “10”.

[codebox]#include <stdio.h>

#include “cutil_inline.h”

device void foo(float *local) {

*local = 10;


global void TestKernel(float *result) {

float local;

local = 5;


*result = local;


int main(int argc, char** argv) {


float *result_d;

cutilSafeCall(cudaMalloc((void**)&result_d, sizeof(float)));

    TestKernel<<<1, 1>>>(result_d);


float result_h;

cutilSafeCall(cudaMemcpy(&result_h, result_d, sizeof(float), cudaMemcpyDeviceToHost));

printf("%g\n", result_h);


return 0;


Compiled with:

[font=“Courier New”]nvcc -gencode arch=compute_20,code=sm_20 -g -G \

            -I~/NVIDIA_GPU_Computing_SDK/C/common/inc/ \

   -o test.cuda[/font]

[codebox]Breakpoint 1, TestKernel <<<(1,1),(1,1,1)>>> (result_d=0x5120000) at

10 local = 5;

(cuda-gdb) s

[Current CUDA Thread <<<(0,0),(0,0,0)>>>]

TestKernel <<<(1,1),(1,1,1)>>> (result_d=0x5120000) at

11 foo(&local);

(cuda-gdb) p local

$3 = 5

(cuda-gdb) s

[Current CUDA Thread <<<(0,0),(0,0,0)>>>]

TestKernel <<<(1,1),(1,1,1)>>> (result_d=0x5120000) at

5 *local = 10;

(cuda-gdb) p local

warning: Variable is not live at this point. Returning garbage value.

$4 = (const @global float * @register) 0x4

(cuda-gdb) s

[Current CUDA Thread <<<(0,0),(0,0,0)>>>]

TestKernel <<<(1,1),(1,1,1)>>> (result_d=0x5120000) at

12 *result_d = local;

(cuda-gdb) p local

warning: Variable is not live at this point. Returning garbage value.

$5 = (const @global float * @register) 0x41200000


[font=“Courier New”]/proc/version

Linux version 2.6.18-194.3.1.el5 ( (gcc version 4.1.2 20080704 (Red Hat 4.1.2-48)) #1 SMP Thu May 13 13:08:30 EDT 2010


NVRM version: NVIDIA UNIX x86_64 Kernel Module 195.36.24 Thu Apr 22 19:10:14 PDT 2010

GCC version: gcc version 4.1.2 20080704 (Red Hat 4.1.2-48)[/font]

CUDA SDK Version 3.0 Release

R195 Driver Release

Videocard: GTX 480

I think that behaviour is expected.

The major complicating factor is that in compiled device code, the function foo doesn’t really exist. There are no subroutines in compiled device code. All device functions get inline expanded by the compiler. So in your example, I doubt there is enough structure in the debugging symbols for the debugger to faithfully replicate what the source says against what the device assembler contains.

I know they are inlined, but that if this is the expected behavior I don’t think the cuda-gdb deserves a 3.0 version number. This severely hampers debugging capabilities for me. Now, if I want to see what actually happens I have to manually inline my functions into one huge function. Which also means I can’t write separate tests for different parts of the code.

If it’s not a bug, I would like to place a feature request. : D

This might be fixed in 3.1. I know I complained about similar behavior with inlined functions not too long ago…