PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

CUDA-x86.

atomic min/max for real data
Goto page Previous  1, 2
 
Post new topic   Reply to topic    PGI User Forum Forum Index -> Accelerator Programming
View previous topic :: View next topic  
Author Message
mkcolg



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

PostPosted: Thu Feb 11, 2010 5:20 pm    Post subject: Reply with quote

Hi Tuan,

For arrays, you want to use the reduction intrinsic "minval". Something like:
Code:
% cat minreduc.cuf

module minutil

  real, device :: dMinval
  real, device :: dMaxval

contains

  attributes(global) subroutine foo (Ad,N)
    use cudafor
    implicit none

    integer, value :: N
    real, device, dimension(N,N) :: Ad
    integer i, j, tx, ty

    tx = threadidx%x
    ty = threadidx%y
    i = (blockidx%x-1)*16 + tx
    j = (blockidx%y-1)*16 + ty

    Ad(i,j) =  (N*(i-1))+(j-1)

    call syncthreads()
    if (i .eq. 1 .and. j .eq. 1 ) then
       dMinval = minval(Ad)
       dMaxval = maxval(Ad)
    endif
    call syncthreads()

  end subroutine foo

  subroutine testmin ()
    use cudafor
    implicit none
    integer :: N = 64
    real, dimension(N,N) :: A
    real, device, dimension(N,N) :: Ad
    real :: minval, maxval
    type(dim3) :: dimGrid, dimBlock
    A=-1
    dMinval = -1
    dimGrid=dim3(N/16,N/16,1)
    dimBlock=dim3(16,16,1)
    Ad=0
    call foo<<<dimGrid,dimBlock>>>(Ad,N)
    A=Ad
    minval = dMinval
    maxval = dMaxval
    print *, minval, maxval
    print *, A(1,1), A(N,N)
  end subroutine testmin

end module minutil

program testme

  use minutil

  call testmin

end program testme

% pgfortran -o minreduc.out minreduc.cuf
% minreduc.out
    0.000000        4095.000
    0.000000        4095.000
Back to top
View user's profile
Tuan



Joined: 11 Jun 2009
Posts: 233

PostPosted: Wed Feb 17, 2010 4:47 pm    Post subject: Reply with quote

mkcolg wrote:
Hi Tuan,

For arrays, you want to use the reduction intrinsic "minval". Something like:
Code:
% cat minreduc.cuf

module minutil

  real, device :: dMinval
  real, device :: dMaxval

contains

  attributes(global) subroutine foo (Ad,N)
    use cudafor
    implicit none

    integer, value :: N
    real, device, dimension(N,N) :: Ad
    integer i, j, tx, ty

    tx = threadidx%x
    ty = threadidx%y
    i = (blockidx%x-1)*16 + tx
    j = (blockidx%y-1)*16 + ty

    Ad(i,j) =  (N*(i-1))+(j-1)

    call syncthreads()
    if (i .eq. 1 .and. j .eq. 1 ) then
       dMinval = minval(Ad)
       dMaxval = maxval(Ad)
    endif
    call syncthreads()

  end subroutine foo

  subroutine testmin ()
    use cudafor
    implicit none
    integer :: N = 64
    real, dimension(N,N) :: A
    real, device, dimension(N,N) :: Ad
    real :: minval, maxval
    type(dim3) :: dimGrid, dimBlock
    A=-1
    dMinval = -1
    dimGrid=dim3(N/16,N/16,1)
    dimBlock=dim3(16,16,1)
    Ad=0
    call foo<<<dimGrid,dimBlock>>>(Ad,N)
    A=Ad
    minval = dMinval
    maxval = dMaxval
    print *, minval, maxval
    print *, A(1,1), A(N,N)
  end subroutine testmin

end module minutil

program testme

  use minutil

  call testmin

end program testme

% pgfortran -o minreduc.out minreduc.cuf
% minreduc.out
    0.000000        4095.000
    0.000000        4095.000


Hi Mat,
Is minval support double precision also, the document just say it supports real, so I'm not sure if this includes double precision?
Please explain me the difference of using min and minval? Does min cannot be used on GPU?

Thanks,
Tuan
Back to top
View user's profile
mkcolg



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

PostPosted: Wed Feb 17, 2010 8:14 pm    Post subject: Reply with quote

Hi Tuan,

Quote:
Is minval support double precision also, the document just say it supports real, so I'm not sure if this includes double precision?
Yes, minval supports double precision. When they doc say 'real', they mean both kinds.

Quote:
Please explain me the difference of using min and minval?


min determines which of two or more scalar values is the minimum values. minval finds the minimum value of an array.

Code:
Does min cannot be used on GPU?

min can be used on the GPU. Though, your code is trying to find the minimum value of an array, hence the use of minval.

- Mat
Back to top
View user's profile
Tuan



Joined: 11 Jun 2009
Posts: 233

PostPosted: Wed Feb 17, 2010 8:34 pm    Post subject: Reply with quote

mkcolg wrote:
Hi Tuan,

Quote:
Is minval support double precision also, the document just say it supports real, so I'm not sure if this includes double precision?
Yes, minval supports double precision. When they doc say 'real', they mean both kinds.

Quote:
Please explain me the difference of using min and minval?


min determines which of two or more scalar values is the minimum values. minval finds the minimum value of an array.

Code:
Does min cannot be used on GPU?

min can be used on the GPU. Though, your code is trying to find the minimum value of an array, hence the use of minval.

- Mat


All clear. Thank Mat.

Tuan
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 Previous  1, 2
Page 2 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