Detecting NaN inside CUDA Fortran Kernel

  • 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.[/code]

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,

% cat test.cuf

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

   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


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
print *, "VAL=", val
call setnan<<<1,1>>>(val)
print *, "ISNAN? ", X

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

Great, thanks a lot for your time!

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?

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.


Thanks Mat!