0

I have a simple kernel, in which I'm allocating some space using malloc, simply as:

__global__ void chainKernel() {
    float* __restrict__ boo = (float*)malloc(sizeof(float));
    *boo = 0;
    *boo = *boo + 100;
    return;
}

If I put a breakpoint on *boo = *boo + 100 I can't see the contents of *boo. Instead I get Operation is not valid due to the current state of the object next to the variable in the debugger window. If I remove the __restrict__ however, the value is shown correctly. Is this normal behavior?

My system: CUDA 5.5.20, Nsight 3.1.0.13141, Windows 7 x64, VS2010, GeForce GTX Titan.

Greg Smith
  • 11,007
  • 2
  • 36
  • 37
user2412789
  • 113
  • 1
  • 14
  • Not sure but the point of `__restrict__` keyword is to avoid *pointer aliasing* and maybe NSight need to create aliases... My guess is that it is a normal behaviour. – Michael M. Aug 20 '13 at 11:13
  • @Michael: so if I want to debug I have to redefine `#define __restrict__`? – user2412789 Aug 20 '13 at 11:15

1 Answers1

3

One of the benefits of __restrict__ is that it allows the compiler to be more aggressive with optimizations. When you have simple code like this that the compiler can completely optimize away, the __restrict__ keyword may help the compiler do just that.

One of the common reasons for not being able to inspect variables in the debugger is due to compiler optimizations, either locally (a variable going out of scope when you weren't expecting it) or globally (a variable that has been completely optimized away).

Note that the definition of the kernel you've shown in this question does nothing useful. Therefore the compiler may be optimizing things away.

To work around this (for this case), put a printf("%f", *boo); statement immediately after the final boo assignment, and the compiler will not be able to optimize the variable away. You should also be using the -G switch for debugging.

Robert Crovella
  • 143,785
  • 11
  • 213
  • 257