Update co_dot exercise (works on Titan). Correct use of c_sizeof and co_reduce binding.

Author: Damian Rouson

src/extensions/opencoarrays.F90

@@ -36,7 +36,7 @@ module opencoarrays
 #elif defined(COMPILER_LACKS_C_SIZEOF_ASSUMED_RANK)
   use iso_c_binding, only : c_int,c_char,c_ptr,c_loc,c_double,c_int32_t,c_bool,c_funloc,c_ptrdiff_t
 #else
-  use iso_c_binding, only : c_int,c_char,c_ptr,c_loc,c_double,c_int32_t,c_bool,c_funloc,c_ptrdiff_t,c_sizeof
+  use iso_c_binding, only : c_int,c_char,c_ptr,c_loc,c_float,c_double,c_int32_t,c_bool,c_funloc,c_ptrdiff_t,c_sizeof
 #endif
   implicit none
 
@@ -240,8 +240,9 @@ subroutine caf_finalize(argc, argv) bind(C,name="_gfortran_caf_finalize")
     ! PREFIX(registernc) (void* mem,size_t mem_size)
 
     subroutine opencoarrays_registernc(mem,mem_size)  bind(C,name="_gfortran_caf_registernc")
+      import :: c_ptr,c_ptrdiff_t,c_int
       type(c_ptr), intent(in), value :: mem
-      integer(c_ptrdiff_t), intent(in) :: mem_size
+      integer(c_ptrdiff_t), value, intent(in) :: mem_size
     end subroutine
 
     ! C function signature from ../mpi/mpi_caf.c:
@@ -268,7 +269,7 @@ subroutine opencoarrays_co_reduce(a, opr, opr_flags, result_image, stat, errmsg,
 #ifdef COMPILER_SUPPORTS_CAF_INTRINSICS
       bind(C,name="_caf_extensions_co_reduce")
 #else
-      bind(C,name="_gfortran_extensions_co_reduce")
+      bind(C,name="_gfortran_caf_co_reduce")
 #endif
       use iso_c_binding, only : c_ptr,c_funptr,c_int,c_char
       type(c_ptr), intent(in), value :: a
@@ -364,8 +365,7 @@ function opencoarrays_this_image(coarray) bind(C,name="_gfortran_caf_this_image"
 
     ! C function signature from ../mpi/mpi_caf.c:
     ! int PREFIX (num_images) (int, int);
-!   function opencoarrays_num_images(coarray,dim_) bind(C,name="_gfortran_caf_num_images") result(num_images_)
-    function opencoarrays_num_images(coarray,dim_) bind(C,name="_caf_extensions_num_images") result(num_images_)
+    function opencoarrays_num_images(coarray,dim_) bind(C,name="_gfortran_caf_num_images") result(num_images_)
       import :: c_int
       integer(c_int), value, intent(in) :: coarray,dim_
       integer(c_int) :: num_images_
@@ -570,8 +570,8 @@ function c_sizeof(mold) result(c_size_of_mold)
   ! _________________________________________________________________________________
 
     subroutine accelerate(a)
-      real(c_double), intent(in), contiguous :: a(..)
-      call opencoarrays_registernc(c_loc(a),size(a)*c_sizeof(a))  
+      real(c_float), intent(in), contiguous, target :: a(..)
+      call opencoarrays_registernc(c_loc(a),c_sizeof(a))  
     end subroutine
 
     subroutine co_reduce_c_int(a, opr, result_image, stat, errmsg) 
@@ -877,23 +877,27 @@ subroutine co_sum_c_int(a,result_image,stat,errmsg)
   end subroutine
 
   ! Return the image number (MPI rank + 1)
-  function this_image()  result(image_num)
-    use mpi, only : MPI_Comm_rank
-    integer(c_int) :: image_num,ierr
-   !image_num = opencoarrays_this_image(unused)
-    call MPI_Comm_rank(CAF_COMM_WORLD,image_num,ierr) 
-    if (ierr/=0) call error_stop
-    image_num = image_num + 1
-  end function
+   function this_image()  result(image_num)
+#ifndef COMPILER_PROVIDES_MPI
+     use mpi, only : MPI_Comm_rank
+#endif
+     integer(c_int) :: image_num,ierr
+    !image_num = opencoarrays_this_image(unused)
+     call MPI_Comm_rank(CAF_COMM_WORLD,image_num,ierr) 
+     if (ierr/=0) call error_stop
+     image_num = image_num + 1
+   end function
 
   ! Return the total number of images
-  function num_images()  result(num_images_)
-    use mpi, only : MPI_Comm_size
-    integer(c_int) :: num_images_,ierr
-   !num_images_ = opencoarrays_num_images(unused_coarray,unused_scalar)
-    call MPI_Comm_size(CAF_COMM_WORLD,num_images_,ierr) 
-    if (ierr/=0) call error_stop
-  end function
+   function num_images()  result(num_images_)
+#ifndef COMPILER_PROVIDES_MPI
+     use mpi, only : MPI_Comm_size
+#endif
+     integer(c_int) :: num_images_,ierr
+    !num_images_ = opencoarrays_num_images(unused_coarray,unused_scalar)
+     call MPI_Comm_size(CAF_COMM_WORLD,num_images_,ierr) 
+     if (ierr/=0) call error_stop
+   end function
 
   ! Halt the execution of all images
   subroutine error_stop(stop_code)

src/tests/integration/gpu/accelerated_co_dot.f90

@@ -50,7 +50,7 @@ subroutine co_dot_accelerated(x,y,x_dot_y,API)
      integer(c_int), intent(in), optional :: API
      integer(c_int) :: chosen_API
 
-     if (present(API)) 
+     if (present(API)) then
        chosen_API = API
      else
        chosen_API = CUDA
@@ -76,6 +76,7 @@ program cu_dot_test
   implicit none
 
   ! Unaccelerated variables 
+  real(c_float), allocatable :: a(:),b(:)
   real(c_float), allocatable :: a_unacc(:),b_unacc(:)
   real(c_float) :: dot
   real(c_double) :: t_start, t_end
@@ -84,7 +85,7 @@ program cu_dot_test
   real(c_float), allocatable :: a_acc(:)[:], b_acc(:)[:]
   real(c_float) :: dot_acc[*]
 
-  integer(c_int),parameter :: n = 99900000
+  integer(c_int) :: n = 99999904
   integer(c_int) :: n_local,np,me
 
   np = num_images()
@@ -121,11 +122,12 @@ program cu_dot_test
 contains
 
   subroutine initialize_all_variables()
+    use opencoarrays,  only : accelerate
     integer(c_int) :: i
     ! These allocation arguments must be coarrays to support the scatter operation below
-    allocate(a_unacc(n_local)[*],b_unacc(n_local)[*])
+    allocate(a_unacc(n_local),b_unacc(n_local))
     ! These allocation arguments will be defined locally and therefore need not be coarrays
-    allocate(a_acc(n_local),b_acc(n_local))
+    allocate(a_acc(n_local)[*],b_acc(n_local)[*])
 
     ! Register the desired variables for acceleration
     call accelerate(a_acc)
@@ -141,10 +143,10 @@ subroutine initialize_all_variables()
         a_acc(1:n_local)[i] = a(n_local*(i-1)+1:n_local*i)
         b_acc(1:n_local)[i] = b(n_local*(i-1)+1:n_local*i)
       end do
-      sync all
-      a_unacc=a_acc
-      b_unacc=b_acc
     endif
+    sync all
+    a_unacc=a_acc
+    b_unacc=b_acc
   end subroutine
 
 end program

src/tests/integration/gpu/co_dot.f90

@@ -1,12 +1,13 @@
 #include <stdio.h>
 #include <stdlib.h>
 #include <cuda.h>
+#include "cublas_v2.h"
 
 extern "C" void cudaPrint(int *data, int n, int me, char *name);
 extern "C" void cudaAdd(int* data_in1, int* data_in2, int *data_out, int n);
 extern "C" void cudaDot(float* in1, float* in2, float* out, int n);
 
-#define MAX_BLOCK_SZ 512
+#define MAX_BLOCK_SZ 1024
 
 __global__ void printOnCuda(int *data, int n)
 {
@@ -27,11 +28,11 @@ __global__ void assignOnCuda(int *data, int n)
 //      printf("From CUDA data[%d] = %d\n",i,data[i]);
 }
 
-__global__ void Dev_dot(float x[], float y[], float z[], int n) {
+__global__ void Dev_dot(float x[], float y[], float z[], long n) {
    /* Use tmp to store products of vector components in each block */
    /* Can't use variable dimension here                            */
    __shared__ float tmp[MAX_BLOCK_SZ];
-   int t = blockDim.x * blockIdx.x + threadIdx.x;
+   long t = blockDim.x * blockIdx.x + threadIdx.x;
    int loc_t = threadIdx.x;
 
    if (t < n) tmp[loc_t] = x[t]*y[t];
@@ -70,18 +71,28 @@ extern "C"
 void cudaDot(float *in1, float *in2, float *out, int n)
 {
   float *partial_dot;
-  int nThreads = 64, i=0;
+  int nThreads = 1024, i=0;
   int nBlocks = ((n-1)/nThreads)+1;
+  float *d_in1,*d_in2;
+//  cublasHandle_t handle;
 
   cudaMallocManaged(&partial_dot, nBlocks * sizeof(float));
+  cudaDeviceSynchronize();
 
   *out = 0.0;
+  
+  cudaHostGetDevicePointer((void **) &d_in1, (void *) in1, 0);
+  cudaHostGetDevicePointer((void **) &d_in2, (void *) in2, 0);
 
-  Dev_dot<<<nBlocks,nThreads>>>(in1,in2,partial_dot,n);
+  Dev_dot<<<nBlocks,nThreads>>>(d_in1,d_in2,partial_dot,n);
+//  cublasCreate (&handle);
+//  cublasSdot(handle,n,d_in1,1,d_in2,1,out);
   cudaDeviceSynchronize();
 
   for(i=0;i<nBlocks;i++)
-    *out += partial_dot[i];
+    {
+      *out += partial_dot[i];
+    }
 
   cudaFree(partial_dot);