Detecting NaN inside CUDA Fortran Kernel

Accelerated Computing HPC Compilers Legacy PGI Compilers
1
  • the subject basically says everything. How can I detect NaN in a kernel? CUDA appears to have a library function[1], but I can’t figure out how to access it.
(my_value .NE. my_value)

returns False also for NaN value, so that approach doesn’t seem to work.

Some background on this: I may have found a bug that would lead to an unspecified launch failure if and only if an array is passed in after a NaN is being passed by value. I want to include some safety check for NaNs in the debug mode of my kernel code generator, such that these kinds of errors can be found quickly.

http://developer.download.nvidia.com/compute/cuda/4_1/rel/toolkit/docs/online/group__CUDA__MATH__DOUBLE_gf9a5caca44978ea06bfcd2345032e5b8.html[/code]

2

Hi MuellerM,

You can call CUDA C routines by adding an interface. Here’s an example of calling “isnan”. The only difficulty is determining what the CUDA C name is. In this case I had to go look through the CUDA C header files since “isnan” gets preprocessed to the particular routine depending upon data type.

Hope this helps,
Mat

% cat test.cuf

module bar
use cudafor
use iso_c_binding
real(8), dimension(:), allocatable, device :: Xd
real(8), dimension(:), allocatable         :: X

interface
   attributes(device) function isnan(a) bind(C,name="__isnand")
      use iso_c_binding
      integer(c_int) :: isnan
      real(c_double), value :: a
   end function isnan
end interface

contains

attributes(global) subroutine setnan (val)
  use iso_c_binding
  real(8),value :: val
  integer :: idx
  integer(c_int) :: nan
  idx = (blockidx%x-1)*blockdim%x + threadidx%x
  nan = isnan(val)
  Xd(idx) = dble(nan)
end subroutine
end module bar

program foo
use cudafor
use bar
real(8) val
val=sqrt(-1.0)
print *, "VAL=", val
allocate(X(1),Xd(1))
call setnan<<<1,1>>>(val)
X=Xd
print *, "ISNAN? ", X

val=sqrt(10.0)
print *, "VAL=", val
call setnan<<<1,1>>>(val)
X=Xd
print *, "ISNAN? ", X
end program foo
% pgf90 test.cuf; a.out
 VAL=                       NaN
 ISNAN?     1.000000000000000
 VAL=    3.162277698516846
 ISNAN?     0.000000000000000
3

Great, thanks a lot for your time!

4

What if I want to use isnan function in C inside an Openacc kernel? I tried with isnan function from math.h but it is not recognized. So how can i solve it?

5

Hi jcastro9999,

We have an open feature request (TPR#21342) for isnan, as well as isfinite, isinf, isnormal, etc., but haven’t added this support to the compiler. I’ll add your request to this TPR and see about increasing the priority.

-Mat

6

Thanks Mat!