module device_code

contains

! CUDA Kernel !

attributes(global) subroutine eval(a,b,c,size)
  implicit none
  integer, value :: size
  real, dimension(0:size-1) :: a,c
  real, value :: b
  integer :: i
  do i = 0, size-1
    c(i)=a(i)*b;
  end do
end subroutine eval

end module device_code

program exercise01

use cudafor
use device_code

implicit none

! define vector size !
integer, parameter :: vector_size = 10000
real, dimension(0:vector_size-1) :: vector_a_cpu
real, dimension(0:vector_size-1) :: vector_c_cpu
real, dimension(0:vector_size-1), device :: vector_a_gpu
real, dimension(0:vector_size-1), device :: vector_c_gpu
real :: b
integer :: i
logical :: success
integer :: error

! CPU code !

call random_seed()

call random_number(b)

! initialize cpu input vector to 0 !
do i=0,vector_size-1
  call random_number(vector_a_cpu(i))
end do
! initialize cpu output vector to 0 !
do i=0,vector_size-1
  vector_c_cpu(i)=0;
end do

! copy input vector to gpu !
vector_a_gpu=vector_a_cpu
! initialize gpu output vector to 0 !
vector_c_gpu=vector_c_cpu

! cuda kernel call !
call eval<<<1,1>>>(vector_a_gpu,b,vector_c_gpu,vector_size)

! copy output vector from gpu to cpu !
vector_c_cpu=vector_c_gpu

! verify results !
success=.true.
do i=0,vector_size-1
  if (abs(vector_a_cpu(i)*b-vector_c_cpu(i))>1e-5*b) then
    print *,'computation result is wrong: index=',i, &
        ' expected result=',vector_a_cpu(i)*b, &      
        'cuda result=', vector_c_cpu(i)
    success=.false.
    exit
  end if
end do
if (success) then
  print *,'computation result is correct.'
end if

end program exercise01