How to print the global variables in cuda-gdb?

I was trying debugging the Cuda Fortran code using cuda-gdb. However, I got wrong values when I printed the global variables inside cude-gdb.
Here is the source code:

module test
  attributes(global) subroutine saxpy(x,y,a)
    implicit none
    real,device :: x(:), y(:)
    real,value :: a
    integer :: i, n

    n = size(x)
    i = blockDim%x * (blockIdx%x - 1) + threadIdx%x
    if (i <= n) y(i) = y(i) + a*x(i)

  end subroutine saxpy
end module test

program testSaxpy
  use test
  use cudafor
  implicit none
  integer, parameter :: N = 40000
  real :: x(N), y(N), result(N), a
  real, device :: x_d(N), y_d(N)
  type(dim3) :: grid, tBlock
  integer :: i

  tBlock = dim3(256,1,1)
  grid = dim3(ceiling(real(N)/tBlock%x),1,1)

  a = 2.0

  do i = 1,N
    x(i) = float(i)
    y(i) = 3.0*float(i)

  result = a*x + y

  x_d = x
  y_d = y
  call saxpy<<<grid, tBlock>>>(x_d,y_d,a)
  y = y_d
  write(*,*) 'Max error: ', maxval(abs(y-result))
end program testSaxpy

The file was compiled using the following command:
pgf90 -Mcuda=cc8.0 testSaxpy.cuf -O0 -g -Mcuda=debug -Mcuda=nordc
Then I ran:
cuda-gdb ./a.out
Then I set a break point:
b 10
Then run the code:
Then at the break point, I tried to print the value of x(1):
p x(1)
and it displayed:
$1 = -6.75539944e+15
which was not correct.

Did I miss anything?
Please let me know if you know how to print the global variables correctly in cuda-gdb. Thanks a lot!


This is a known issue with support for Fortran metadata in CUDA that required changes in both CUDA and nvfortran (aka pgf90). I tested your code with our development compiler and the issue appears to have been resolved. Assuming it passes our internal testing, the changes should be available in an upcoming release.


Thanks a lot for the information Mat!
When will the new release available? I am looking forward to it.

Unfortunately I can’t give out release dates but we put out new releases every few months, so don’t expect it to be too long. Though again this is with the caveat that it still needs to pass our QA so it’s not guaranteed to be in the next release.