PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

Free OpenACC Webinar

Finding minval and maxval using CUF kernels
Goto page Previous  1, 2, 3
 
Post new topic   Reply to topic    PGI User Forum Forum Index -> Programming and Compiling
View previous topic :: View next topic  
Author Message
mkcolg



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

PostPosted: Thu Mar 06, 2014 2:04 pm    Post subject: Reply with quote

Quote:
what version of PGF do I need for atomicinc? I am using 13.10. It doesnt seem to be incrementing NodesInCell as hoped.
We just add them to OpenACC in 14.1, but they've been in CUDA Fortran for quite awhile. So my example wont work for you but you should be able to write you're own CUDA Fortran kernel and use them there.

- Mat
Back to top
View user's profile
SPHriction-3D



Joined: 12 Jan 2014
Posts: 48

PostPosted: Thu Mar 06, 2014 5:25 pm    Post subject: Reply with quote

I tried this, but NodesInCell is not being incremented. Each entry of NodesInCell is zero.

Code:
Attributes(Global) Subroutine BuildCellListAtomic(nTotal,CellIndex,CellList,maxIndex,NodesInCell)

   Implicit None

   Integer:: i, j
   Integer, Device:: idx
   Integer, Value:: nTotal, maxIndex   
   Integer, Device, Intent(IN):: CellIndex(nTotal,2)
   Integer, Device, Intent(OUT):: CellList(maxIndex,27), NodesInCell(maxIndex)

   CellList      = 0
   NodesInCell   = 0

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

   If (i >= 1 .and. i <= nTotal) Then
      j = CellIndex(i,2)   
      idx = atomicinc(NodesInCell(j),27)
      idx = idx + 2         
      CellList(j,idx) = CellIndex(i,1)      
      CellList(j,1) = j   
   End If 
         
End Subroutine BuildCellListAtomic


Have I done something wrong?

Kirk
Back to top
View user's profile
mkcolg



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

PostPosted: Fri Mar 07, 2014 10:01 am    Post subject: Reply with quote

You shouldn't clear "NodesInCell" in the kernel itself. This will get run by all threads, in any order, and will most likely give you zeros in the result.

Here's a small example:

Code:
% cat test_3_7_14.cuf
module cells

contains

attributes(Global) Subroutine BuildCellListAtomic(nTotal,CellIndex,CellList,maxIndex,NodesInCell)

    Implicit None

    Integer:: i, j
    Integer, Device:: idx
    Integer, Value:: nTotal, maxIndex
    Integer, Device, Intent(IN):: CellIndex(nTotal,2)
    Integer, Device, Intent(OUT):: CellList(maxIndex,27), NodesInCell(maxIndex)

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

    If (i >= 1 .and. i <= nTotal) Then
       j = CellIndex(i,2)
       idx = atomicinc(NodesInCell(j),27)
       idx = idx + 2
       CellList(j,idx) = CellIndex(i,1)
       CellList(j,1) = j
    End If

 End Subroutine BuildCellListAtomic


end module cells


Program Bounds

 Use cudafor
 use cells
 Implicit None

    Integer, parameter :: nTotal=1024
    Integer, parameter :: maxIndex=64
    Integer :: CellIndex(nTotal,2)
    Integer :: CellList(maxIndex,27), NodesInCell(maxIndex)
    Integer, Device :: CellIndex_d(nTotal,2)
    Integer, Device :: CellList_d(maxIndex,27), NodesInCell_d(maxIndex)
    Integer, Parameter:: BlockX = 128, BlockY = 1
    Integer :: i
    type(dim3) :: tBlock, grid

    NodesInCell=0
    CellIndex=0
    CellList=0
    do i=1,nTotal
        CellIndex(i,2) = mod(i,maxIndex-1)+1
        CellIndex(i,1) = i
    enddo
    NodesInCell_d=NodesInCell
    CellIndex_d=CellIndex
    CellList_d=CellList
    tBlock = Dim3(BlockX,BlockY,1)
    grid = Dim3(ceiling(Real(nTotal)/tBlock%x),1,1)
    Call BuildCellListAtomic<<<grid,tBlock>>>(nTotal,CellIndex_d,CellList_d,maxIndex,NodesInCell_d)
    NodesInCell=NodesInCell_d
    CellList=CellList_d
    print *, NodesInCell
 End Program Bounds

% pgf90 test_3_7_14.cuf -V13.10; a.out
           16           17           17           17           17           17
           17           17           17           17           17           17
           17           17           17           17           17           16
           16           16           16           16           16           16
           16           16           16           16           16           16
           16           16           16           16           16           16
           16           16           16           16           16           16
           16           16           16           16           16           16
           16           16           16           16           16           16
           16           16           16           16           16           16
           16           16           16            0
Back to top
View user's profile
SPHriction-3D



Joined: 12 Jan 2014
Posts: 48

PostPosted: Sat Mar 08, 2014 1:16 pm    Post subject: Reply with quote

Fantastic!

That is perfect, makes a big difference.

Thank you, your help is very appreciated.

Kirk
Back to top
View user's profile
Display posts from previous:   
Post new topic   Reply to topic    PGI User Forum Forum Index -> Programming and Compiling All times are GMT - 7 Hours
Goto page Previous  1, 2, 3
Page 3 of 3

 
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