PGF90-S-0155-Calls from device code to a host function are a

The following program compiles successfully:
module mod
contains
attributes(global) subroutine one(a,nthread,nblock)
implicit none
integer, value :: nthread,nblock
integer, device :: a(nthread,nblock)
! integer :: two
a(threadidx%x,blockidx%x)=two(a(threadidx%x,blockidx%x))
a(threadidx%x,blockidx%x)=two(a(threadidx%x,blockidx%x))
end subroutine one
attributes(device) integer function two(a)
implicit none
integer :: a
two=a+threadidx%x+10blockidx%x
end function two
end module mod
program main
use mod
integer, device, allocatable :: d_a(:,:),d_nthread,d_nblock
integer, allocatable :: a(:,:)
nthread=2
nblock=2
allocate(d_a(nthread,nblock))
allocate(a(nthread,nblock))
d_a=0
call one<<<nthread,nblock>>>(d_a,nthread,nblock)
istat=cudathreadsynchronize()
a=d_a
write(
,*) a
end program main

However, if I uncomment the line
! integer :: two
in subroutine one, I get the following compiler messages:
PGF90-S-0155-Calls from device code to a host function are allowed only in emulation mode - two (tst.cuf: 8)
PGF90-S-0155-Calls from device code to a host function are allowed only in emulation mode - two (tst.cuf: 9)

What the compiler is trying to tell me?

Peter Nightingale

Hi Peter Nightingale,

When you declare “integer :: two”, you’re actually declaring an external function. So when the compiler reaches the use of “two” it issues the syntax error since at this point, it doesn’t know that you really mean for the external “two” to be the same as the module’s “two” routine.

I have submitted a technical problem report (TPR#17672) to see if we can delay this check until after the entire module is processed and all internal reference have be resolved.

Thanks,
Mat

Some late housekeeping. This TPR has been fixed as of the 12.8 release.

thanks,
dave

Hi.

I am having the same compilation error once I try to compile the following (extremely basic) kernel. The kernel stands alone in a module file called kern_strided.cuf

module kern_strided

  implicit none
  contains

  attributes(global) subroutine stride(a, jump)
    real, intent(inout) :: a
    integer, intent(in), value :: jump

    integer :: i, is

    i  = blockDim%x * (blockIdx%x - 1) + threadIdx%x
    is = i * jump
    a(is) = a(i) + 1

  end subroutine stride

end module kern_strided

And I compile it with PGI/17.4 like:

pgfortran -Mcuda=cc50,cuda8.0 -c kern_strided.cuf -o kern_strided.o

with the following error message:

PGF90-S-0155-Calls from device code to a host function are allowed only in emulation mode - a (kern_strided.cuf: 15)
PGF90-S-0155-Calls from device code to a host function are allowed only in emulation mode - a (kern_strided.cuf: 15)
PGF90-S-0072-Assignment operation illegal to external procedure a (kern_strided.cuf: 15)
  0 inform,   0 warnings,   3 severes, 0 fatal for stride

Line 15 is the one where ‘a(is) = a(i) + 1’. Surprisingly, a similar kernel (with similar variable declaration and module construct) compiles successfully, but this one fails to.

Any comment/idea is welcome.

Regards,
Ehsan.

Hi Ehsan,

You have an error in your code. “a” is declared as a scalar and hence “a(is)” is interrupted as a function call. To fix, make “a” an array.

module kern_strided

   implicit none
   contains

   attributes(global) subroutine stride(a, jump)
     real, intent(inout), dimension(:) :: a
     integer, intent(in), value :: jump

     integer :: i, is

     i  = blockDim%x * (blockIdx%x - 1) + threadIdx%x
     is = i * jump
     a(is) = a(i) + 1

   end subroutine stride

 end module kern_strided

Hope this helps,
Mat