[SOLVED] 'pow' function doesn't works with __constant__ variable in debug mode.

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

I was testing my code and I’ve seen that the function ‘pow’ doesn’t work with a constant variable in debug mode.

__constant__ double v;

__global__ void myKernel(...)
{
    ...
    double x = pow(v, 2);
    ...
}

If I debug the code step by step with cuda-gdb the value of ‘x’ is 0 and if I try to print ‘pow(v, 2)’ in cuda-gdb command line I get the message “This target does not support functions calls.”.

This cause that the result of my kernel is not correct, but if I execute the same code without debug mode the result is correct and all works fine.

I don’t understand this behavior, does anyone have any idea why it might be? I understand it is by the constant variable but I don’t know the reason.

Thank you.

2

Many of CUDA’s device-side math functions are inlined, some (especially complex ones, like pow()), are called subroutines, or make use of called subroutines. In addition, function inlining may largely be disabled in debug builds, increasing the number of called functions.

So if the target does not support function calls within some parts of the debugger, you may not be able to single-step through functions implemented as called subroutines. At least that is how I understand the error message, which admittedly I have never encountered before. It also implies the issue is not with the constant data per se, but with the way functions are invoked.

I assume by “target” the debugger refers to the GPU architecture, or more specifically the compute capability. What GPU are you using, on what OS platform? What nvcc compiler switches did you use to build the code? I wonder whether the error message displayed could also be a consequence of compiling with the switch -abi=no.

3

Hi,

I’m using Ubuntu Server 16.04 and GPU Quadro M4000. Capacity is 5.2.

Well, it’s a strange case, I guess the problem is how you said it. I’ll try to run it in another workstation to see if the problem is reproduced too.

I’ll indicate the results.

Thank you

4

I didn’t have any difficulty debugging using a simple test case built around the code you have shown. Here is my full session:

$ cat t71.cu
#include <stdio.h>
#include <math.h>

__constant__ double v;

__global__ void kernel(){

  double x = pow(v, 2);
  printf("x = %f\n", x);
}

int main(){

  double val = 2.0;
  cudaMemcpyToSymbol(v, &val, sizeof(double));
  kernel<<<1,1>>>();
  cudaDeviceSynchronize();
}
$ nvcc -arch=sm_61 -o t71 t71.cu
$ cuda-memcheck ./t71
========= CUDA-MEMCHECK
x = 4.000000
========= ERROR SUMMARY: 0 errors
$ nvcc -g -G -arch=sm_61 -o t71 t71.cu
$ cuda-memcheck ./t71
========= CUDA-MEMCHECK
x = 4.000000
========= ERROR SUMMARY: 0 errors
$ cuda-gdb ./t71
NVIDIA (R) CUDA Debugger
8.0 release
Portions Copyright (C) 2007-2016 NVIDIA Corporation
GNU gdb (GDB) 7.6.2
Copyright (C) 2013 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-unknown-linux-gnu".
For bug reporting instructions, please see:
<http://www.gnu.org/software/gdb/bugs/>...
Reading symbols from /home/nvidia/xxxx/t71...done.
(cuda-gdb) break t71.cu:9
Breakpoint 1 at 0x402b3f: file t71.cu, line 9.
(cuda-gdb) run
Starting program: /home/nvidia/xxxx/./t71
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
[New Thread 0x7ffff04ba700 (LWP 7203)]
[New Thread 0x7fffefc38700 (LWP 7204)]
[New Thread 0x7fffef437700 (LWP 7205)]
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]

Breakpoint 1, kernel<<<(1,1,1),(1,1,1)>>> () at t71.cu:9
9         double x = pow(v, 2);
(cuda-gdb) print x
$1 = 7.288035599595993e-304
(cuda-gdb) print v
$2 = 2
(cuda-gdb) next
10        printf("x = %f\n", x);
(cuda-gdb) print x
$3 = 4
(cuda-gdb) print v
$4 = 2
(cuda-gdb)

CUDA 8, GTX Titan X Pascal, Ubuntu 14.04

I think if you issue the command

print pow(v, 2)

in cuda-gdb, you will get that message:

This target does not support function calls.

I would simply interpret that as “This capability is not implemented in cuda-gdb at this time.”

I’m not really sure that gdb implements that either. For example, consider the following gdb session which has essentially nothing to do with cuda:

$ cat t71.cu
#include <stdio.h>
#include <math.h>

__constant__ double v;

__global__ void kernel(){

  double x = pow(v, 2);
  printf("x = %f\n", x);
}

int main(){

  double val = 2.0;
  double x = pow(val, 2);
  cudaMemcpyToSymbol(v, &val, sizeof(double));
  kernel<<<1,1>>>();
  cudaDeviceSynchronize();
}
$ nvcc -g -G -arch=sm_61 -o t71 t71.cu
$ gdb ./t71
GNU gdb (Ubuntu 7.7.1-0ubuntu5~14.04.2) 7.7.1
Copyright (C) 2014 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-linux-gnu".
Type "show configuration" for configuration details.
For bug reporting instructions, please see:
<http://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 ./t71...done.
(gdb) break t71.cu:16
Breakpoint 1 at 0x402aaf: file t71.cu, line 16.
(gdb) run
Starting program: /home/nvidia/xxxx/t71
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".

Breakpoint 1, main () at t71.cu:16
16        double val = 2.0;
(gdb) print val
$1 = 6.9533558074911346e-310
(gdb) next
17        double x = pow(val, 2);
(gdb) print val
$2 = 2
(gdb) print pow(val, 2)
No symbol "pow" in current context.
(gdb)

I’m not a gdb expert, but its not obvious to me that

print pow(...)

is actually supported in gdb

5

Hi,

I tried this issue in other workstation and it worked fine. Also I tried to compile the software in other workstation and executed it in the original workstation and it worked fine too.

I think that it was a problem with nvcc.

thank you.