Problem with PGI 10.2

I have a very simple test code that worked before, but will no longer compile with v10.2 of the pgf90.

module gpudata
integer, device :: d_int(4)
end module gpudata

module gpufunc
contains
attributes(global) subroutine myfunc()
use gpudata
d_int(threadidx%x) = threadidx%x
end subroutine myfunc
end module gpufunc

program fcuda
use gpudata
use gpufunc

implicit none
integer :: h_int(4)

d_int = 0
h_int = d_int
print *, h_int
call myfunc<<<1,4>>>()
h_int = d_int
print *, h_int

end program fcuda

Is this a known bug? And is there a workaround?

Thanks,
Todd

Hi Todd,

A device subprogram can only access a device variable declared within the same module (See section 2.1.1 of the CUDA Fortran programming guide) The reason for this is that there isn’t a linker (at least not yet) for the GPU so there is no way to associate the global variables instantiated in one module with one instantiated in another. In other words, each time you used “gpudata”, a new copy of d_int would be created.

The actual problem is that we’re not catching this semantic error. It just happened to work in this particular case for 10.1. If you moved “gpudata” to a separate file, then you would get the same error as you see now. Our engineers are working on catching this error.

To fix, move “d_init” into “gpufunc” module.

% cat todd.cuf
module gpufunc
integer, device :: d_int(4)
contains
attributes(global) subroutine myfunc()
d_int(threadidx%x) = threadidx%x
end subroutine myfunc
end module gpufunc

program fcuda
use gpufunc

implicit none
integer :: h_int(4)

d_int = 0
h_int = d_int
print *, h_int
call myfunc<<<1,4>>>()
h_int = d_int
print *, h_int

end program fcuda

xps730:/tmp/qa/loop% pgf90 todd.cuf -V10.2
xps730:/tmp/qa/loop% a.out
            0            0            0            0
            1            2            3            4

Hope this helps,
Mat

Hey Mat,

Thanks for the information. In our code, I was carefully keeping our CUDA modules in the same file to avoid problems. I’ve updated the code as you suggested and everything works now.

Obviously, having a linker is a desirable feature — any idea when that might happen?

Thanks,
Todd

Obviously, having a linker is a desirable feature — any idea when that might happen?

Nope, sorry.