CUDAFortan error: ... from device code to a host function...

Dear All, I am learning CUDA Fortran. I want to write a program to calculate the sumations E of n=10 values of random number. Here is my code:

!%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
MODULE sum_random_number
IMPLICIT NONE

CONTAINS
!!!=================================================================
ATTRIBUTES(GLOBAL) SUBROUTINE average_thermal(E)
IMPLICIT NONE

REAL (KIND=8),INTENT(INOUT):: E(:)
INTEGER (KIND=8) :: iT,i
REAL (KIND=8) :: rdn
INTEGER (KIND=8),PARAMETER :: n=10,N_rdn=200

!REAL (KIND=8),device :: deviceData (N_rdn)

iT=threadIdx%x

E(iT)=0.
DO i=1,n
E(iT)=E(iT)+deviceData(i)
END DO

END SUBROUTINE average_thermal

END MODULE sum_random_number

!!!=================================================================
PROGRAM main_sum
USE cudafor
USE sum_random_number
USE curand_m

IMPLICIT NONE
INTEGER (KIND=8),PARAMETER :: nT=15
REAL (KIND=8),DEVICE :: Ed(nT)
REAL (KIND=8) :: E(nT)
INTEGER (KIND=8) :: iT

INTEGER (KIND=8),PARAMETER :: N_rdn=200

REAL (KIND=8),device :: deviceData(N_rdn)
INTEGER (KIND=8) :: gen,seed

CALL curandCreateGenerator(gen,CURAND_RNG_PSEUDO_DEFAULT)
seed=1234
CALL curandSetPseudoRandomGeneratorSeed(gen,seed)
CALL curandGenerateUniform (gen,deviceData, N_rdn)

Ed=E !!! Host to Divice transpers
CALL average_thermal<<<1,nT>>>(Ed)
E=Ed !! Device to Host transpers

OPEN(unit=20,file=‘average_thermal.dat’)
DO iT=1,nT
WRITE(20,) iT,E(iT)
WRITE(
,*) iT,E(iT)
END DO

CLOSE(20)

! Destroy the generator
CALL curandDestroyGenerator(gen)

END PROGRAM main_sum
!%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%

When I compiled, there was the erros: PGF90-S-0155-Calls from device code to a host function are allowed only in emulation mode - devicedata

Please help me to figure out this problem! Thank you in avance.

Hi Danh-Tai HOANG,

Notice the second error message you are getting:

PGF90-S-0038-Symbol, devicedata, has not been explicitly declared

Since “deviceData” isn’t declared in average_thermal’s scope, implicit typing is used and “deviceData(I)” defaults to being typed as an external function call. To fix, you need to me deviceData’s declaration into the module’s data section so it’s visible to both the main program and the subroutine.

  • Mat

Hi mkcolg,
Thank you for your help.
In fact, when “deviceData” is declared in “average_thermal”, there is not error but the values of E are always 0.

Based on your idea, I added “deviceData” as the second input:

average_thermal(E,deviceData)
REAL (KIND=8),device,INTENT(IN) :: deviceData (N_rdn)

And in main program: CALL average_thermal<<<1,nT>>>(Ed,deviceData)

It works now!