Skip to content

Commit

Permalink
new file: inc.cuf
Browse files Browse the repository at this point in the history
	increment kernel
  • Loading branch information
Qengine committed Mar 18, 2022
1 parent d6f55f6 commit 003daaa
Showing 1 changed file with 57 additions and 0 deletions.
57 changes: 57 additions & 0 deletions inc.cuf
Original file line number Diff line number Diff line change
@@ -0,0 +1,57 @@
module simpleOps_m
contains
attributes (global) subroutine increment (a, b)
implicit none
integer , intent (inout) :: a (:)
integer , value :: b
integer :: i , n

i = blockDim%x * (blockIdx%x-1) + threadIdx%x
n = size(a)
if ( i <= n ) a(i) = a(i) + b

end subroutine increment
end module simpleOps_m

program incrementTest
use cudafor
use simpleOps_m
implicit none
integer, parameter :: n = 1024*1024*1024
integer, allocatable :: a(:)
integer, device , allocatable :: a_d(:)
integer :: b , tPB = 256
type(cudaEvent) :: startEvent, stopEvent
integer :: istat
real :: time

allocate(a(n),a_d(n))
a = 1
b = 3

istat = cudaEventCreate(startEvent)
istat = cudaEventCreate(stopEvent)

a_d = a
istat = cudaEventRecord(startEvent, 0)
call increment <<< ceiling(real(n)/ tPB) , tPB >>> (a_d, b)
istat = cudaEventRecord(stopEvent, 0)
istat = cudaEventSynchronize(stopEvent)
istat = cudaEventElapsedTime(time, startEvent, stopEvent)
a = a_d

if (any ( a /= 4)) then
write (* ,*) "**** ProgramFailed ******"
else
write (* ,*) "ProgramPassed "
endif

write(*,*) "Device. kernel took", real(time*1.d-6,4), "sec"

deallocate(a, a_d)

istat = cudaEventDestroy(startEvent)
istat = cudaEventDestroy(stopEvent)

end program incrementTest

0 comments on commit 003daaa

Please sign in to comment.