PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

Free OpenACC Webinar

Module device array as arg to kernel: Should this work?

 
Post new topic   Reply to topic    PGI User Forum Forum Index -> Accelerator Programming
View previous topic :: View next topic  
Author Message
Robert Link



Joined: 14 Jun 2010
Posts: 9
Location: Goddard Space Flight Center

PostPosted: Fri Jul 30, 2010 7:22 am    Post subject: Module device array as arg to kernel: Should this work? Reply with quote

Consider the following kernel:
Code:

  attributes(global) subroutine reorder(ncs, N1, N2, jmax, cout, cin)
    integer, value :: ncs, N1, N2, jmax
    real,device :: cout(N1,N2), cin(N1, N2)

    integer :: tid, bn, indx
    integer :: jnew, iout

   
    tid   = threadidx%x
    bn    = (blockidx%y-1) * griddim%x + blockidx%x-1
    indx  = bn*blockdim%x + tid

       iout  = jreorder(indx)
       
       if(iout .ge. 1 .and. iout .le. N1) then
         
          do jnew=1,jmax
             cout(iout,jnew) = cin(indx,jnew)
          end do
       end if
   
  end subroutine reorder

where jreorder is a permutation vector (i.e., guaranteed to have no duplicate entries) that is declared in the enclosing module. The kernel is called thus:
Code:

call unreorder<<<gsize,bsize>>>(ncs, NCELL, MXGSAER, ischang(ncs), cx, cnew)

where, as it happens, cx and cnew are also declared in the enclosing module and are declared (or allocated -- I've tried both) as NCELL,MXGSAER. Also, the 4th argument is verified to be less than MXGSAER, so we are guaranteed not to exceed the array bounds.

So far as I can tell, this should work. However, when I run the code as described, the results are incorrect, and in fact they change when you switch from statically declared to run-time allocated arrays (suggesting that faulty memory access is to blame). If, on the other hand, I change the kernel to this:
Code:

  attributes(global) subroutine reordercx(ncs,jmax)
    integer, value :: ncs,ischan

    integer :: tid, bn, indx
    integer :: jnew,iout

    tid   = threadidx%x
    bn    = (blockidx%y-1) * griddim%x + blockidx%x-1
    indx  = bn*blockdim%x + tid

       iout = jreorder(indx)
       do jnew=1,jmax
          cx(iout, jnew) = cnew(indx,jnew)
       end do
  end subroutine reordercx

taking advantage of the fact that cx and cnew are available as module variables, the answers are correct for both the declared and allocated versions.

Is this a compiler bug, or am I overlooking some error I committed in the first version of the kernel?

-robert.
Back to top
View user's profile
mkcolg



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

PostPosted: Fri Jul 30, 2010 3:19 pm    Post subject: Reply with quote

Hi Robert,

This is known issue that we're actively working on. I've talked to our Compiler Engineering Manager and this is one of their high priority items for September's 10.9 release. While, we can't guarantee it will be fixed in 10.9, it is the goal.

FYI, the problem report number is TPR#16790.

Thanks,
Mat
Back to top
View user's profile
Robert Link



Joined: 14 Jun 2010
Posts: 9
Location: Goddard Space Flight Center

PostPosted: Sat Jul 31, 2010 9:25 am    Post subject: Reply with quote

mkcolg wrote:
Hi Robert,

This is known issue that we're actively working on. I've talked to our Compiler Engineering Manager and this is one of their high priority items for September's 10.9 release. While, we can't guarantee it will be fixed in 10.9, it is the goal.



Thanks for the info, Mat. Is the bug limited to passing module arrays as arguments to kernels, or are all array arguments affected?

-robert.
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
Page 1 of 1

 
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