PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

Free OpenACC Webinar

cuda variable not updating
Goto page 1, 2  Next
 
Post new topic   Reply to topic    PGI User Forum Forum Index -> Accelerator Programming
View previous topic :: View next topic  
Author Message
cablesb



Joined: 21 Jan 2010
Posts: 33

PostPosted: Mon May 21, 2012 11:26 am    Post subject: cuda variable not updating Reply with quote

I am trying to use CUDA to run what I thought would be a very simple-minded coding of the diffusion equation. Here's the code:
Code:

module dimensions

integer, parameter :: nx=256, ny=256

end module




program diffuse

use cudafor
use dimensions

real*8 :: v(nx,ny), diffconst
real*8, device :: v_d(nx,ny)
integer :: nloops, outputdl
type(dim3) :: grid, tBlock
integer, device :: n_d, outputdl_d
real*8, device :: diffconst_d

nloops=10000
outputdl=1000
diffconst=.25

outputdl_d=outputdl
diffconst_d=diffconst

tBlock=dim3(64,64,1)
grid=dim3(ceiling(real(nx)/tBlock%x), &
          ceiling(real(ny)/tBlock%y), 1)

open(unit=11,file='diff_output.txt')

v=0

n=0

do while (n.le.nloops)
  v_d=v
  n_d=n
  call diff_time_stepper<<<grid>>>(v_d,diffconst_d,n_d,outputdl_d)
  v=v_d
  n=n_d
  write(11,*) v
  print *,n
enddo


end program





attributes(global) subroutine diff_time_stepper(v,diffconst,n,nsteps)


real*8 :: v(:,:)
real*8 :: diffconst
integer :: n
integer :: nsteps

real*8 :: vintermed
integer :: i,j,m
integer :: nx, ny

nx=256
ny=256

i=(blockIdx%x-1)*blockDim%x+threadIdx%x
j=(blockIdx%y-1)*blockDim%y+threadIdx%y

do m=1,nsteps
  if (i<nx .and. j<ny>1 .and. j>1) then
    vintermed=v(i,j)+diffconst*(v(i-1,j)-2.*v(i,j)+v(i+1,j)+v(i,j-1)-2.*v(i,j)+v(i,j+1))
    v(i,j)=vintermed
  endif
! add a source for the heck of it
  if (i==64 .and. j==64) v(i,j)=v(i,j)+1
enddo


n=n+nsteps

end subroutine


The one obvious problem I have is is that the variable "n" (returned to main as "n_d") is not updating. At all. There are probably numerous other problems i don't know about too. Trying to run with cuda=emu gives a seg fault. Any ideas? THanks. I'm really new to this.
Back to top
View user's profile
cablesb



Joined: 21 Jan 2010
Posts: 33

PostPosted: Mon May 21, 2012 11:28 am    Post subject: Reply with quote

But wait, there's more. In the triple chevron, I don't know what happened. The code -- I am looking at it now -- reads

call diff_time_stepper<<<grid>>>( ... )
Back to top
View user's profile
cablesb



Joined: 21 Jan 2010
Posts: 33

PostPosted: Mon May 21, 2012 11:29 am    Post subject: Reply with quote

OK, this is like a Twilight Zone episode or something. Inside the triple chevrons are "grid,tBlock". Don't know how to make it look right.
Back to top
View user's profile
TheMatt



Joined: 06 Jul 2009
Posts: 322
Location: Greenbelt, MD

PostPosted: Mon May 21, 2012 11:59 am    Post subject: Reply with quote

cablesb wrote:
OK, this is like a Twilight Zone episode or something. Inside the triple chevrons are "grid,tBlock". Don't know how to make it look right.

I think this has to do with the HTML parser or something with the Forum. Try substituting, say, {{{ }}} instead in your paste.
Back to top
View user's profile
mkcolg



Joined: 30 Jun 2004
Posts: 6211
Location: The Portland Group Inc.

PostPosted: Mon May 21, 2012 1:14 pm    Post subject: Reply with quote

Hi cablesb,

Quote:
I think this has to do with the HTML parser or something with the Forum. Try substituting, say, {{{ }}} instead in your paste.
Yes, it's our UF. I keep meaning on looking at the UF code to see if can fix it, but just haven't done it yet.

Quote:
The one obvious problem I have is is that the variable "n" (returned to main as "n_d") is not updating. At all. There are probably numerous other problems i don't know about too. Trying to run with cuda=emu gives a seg fault. Any ideas? THanks. I'm really new to this.
Your code has a number of issues. First is that global kernel routines *must* have an interface. Without an interface, your program will most likely seg fault if you are passing arguments. You can fix this by adding an explicit interface, or put your routine in a module. Module routines are given an implicit interface.

Next, your algorithm will give you wrong answers in two spots:
Code:
n=n+nsteps
Here every thread will update "n" and the value stored back to global memory will be dependent upon the execution order of the threads. To fix, you need to either guard the updates (like "if (i.eq.1.and.j.eq.1) then") or better yet, move this to the host code.
Code:
 
do m=1,nsteps
  if (i<nx .and. j<ny>1 .and. j>1) then
    vintermed=v(i,j)+diffconst*(v(i-1,j)-2.*v(i,j)+v(i+1,j)+v(i,j-1)-2.*v(i,j)+v(i,j+1))
    v(i,j)=vintermed
  endif
! add a source for the heck of it
  if (i==64 .and. j==64) v(i,j)=v(i,j)+1
enddo

Again, this code will lead to non-deterministic results. You update the new value v(i,j) before making sure all threads that need the old value have read it. Hence, some threads may be reading the new value, and others the old. Worse, the order in which the threads changes from run to run and so will your results.

You need to globally synchronize your threads in between the reads and write to "v" and the only method to globally synchronize threads is between kernel calls. Hence, you need to change your algorithm to something like:
Code:

! Host code
do while (n.le.nloops)
    do m=1,nsteps
!         call kernel1 to perform the calculation and store the results to a temp array
!         copy temp array back into v
    enddo
    n=n+nsteps
end do


I use a similar algorithm for my CUDA Fortran implementation of the "Game of Life" (found HERE).

Hope this helps,
Mat
Back to top
View user's profile
Display posts from previous:   
Post new topic   Reply to topic    PGI User Forum Forum Index -> Accelerator Programming All times are GMT - 7 Hours
Goto page 1, 2  Next
Page 1 of 2

 
Jump to:  
You cannot post new topics in this forum
You cannot reply to topics in this forum
You cannot edit your posts in this forum
You cannot delete your posts in this forum
You cannot vote in polls in this forum


Powered by phpBB © phpBB Group