Adding cuada_mpi and co_dot integration test.

Author: Damian Rouson

src/cuda_mpi/mpi_caf.c

@@ -364,20 +364,33 @@ PREFIX (num_images)(int distance __attribute__ ((unused)),
   return caf_num_images;
 }
 
+/* Register user-allocated memory for management by CUDA (works for coarrays and non-coarrays):
+   This procedure is exposed in the OpenCoarrays Fortran application programming interface (API)
+   contained in ../extensions/opencoarrays.f90 to expose it to Fortran programmers who want
+   to explicitly allocate memory related to manycore devices: currently via CUDA for NVIDIA
+   GPUs and possibly in the future via another mechanism for Intel MIC architecture processors.
+*/
+
 void 
-PREFIX(registernc) (size_t size,void* mem)
+PREFIX(registernc) (void* mem,size_t mem_size)
 {
   int cuda_ierr = 0;
 
-  cuda_ierr = cudaHostRegister(mem,size,cudaHostRegisterMapped);
+  cuda_ierr = cudaHostRegister(mem,mem_size,cudaHostRegisterMapped);
   cudaDeviceSynchronize();
 
-  if (ierr != 0) call caf_runtime_error ("CUDA allocation failed with code %d", ierr);
+  if (ierr != 0) 
+    call caf_runtime_error ("CUDA allocation failed with code %d", ierr);
 
   return;
 
 }
 
+/* Allocate coarray variables:
+   This procedure is part of the OpenCoarrays Fortran application binary interface (ABI)
+   that is used by the compiler only and is therefore not exposed to Fortran programs directly.
+*/
+
 void *
 PREFIX (register) (size_t size, caf_register_t type, caf_token_t *token,
                   int *stat, char *errmsg, int errmsg_len, int cuda)

src/extensions/opencoarrays.F90

@@ -52,6 +52,7 @@ module opencoarrays
   public :: sync_all
   public :: caf_init
   public :: caf_finalize
+  public :: accelerate
 #ifdef COMPILER_SUPPORTS_ATOMICS
   public :: event_type
   public :: event_post
@@ -216,6 +217,7 @@ pure function logical_operator(lhs,rhs) result(lhs_op_rhs)
 
   ! Bindings for OpenCoarrays C procedures
   interface 
+
     ! C function signature from ../mpi/mpi_caf.c
     ! void _gfortran_caf_init (int *argc, char ***argv);
 
@@ -233,6 +235,15 @@ subroutine caf_finalize(argc, argv) bind(C,name="_gfortran_caf_finalize")
       type(c_ptr), value ::  argv
     end subroutine
 
+    ! C function signature from ../cuda_mpi/mpi_caf.c:
+    ! void 
+    ! PREFIX(registernc) (void* mem,size_t mem_size)
+
+    subroutine opencoarrays_registernc(mem,mem_size)  bind(C,name="_gfortran_caf_registernc")
+      type(c_ptr), intent(in), value :: mem
+      integer(c_ptrdiff_t), intent(in) :: mem_size
+    end subroutine
+
     ! C function signature from ../mpi/mpi_caf.c:
     ! void
     ! PREFIX (co_min) (gfc_descriptor_t *a, int result_image, int *stat, char *errmsg,
@@ -558,6 +569,11 @@ function c_sizeof(mold) result(c_size_of_mold)
   ! ______ Assumed-rank co_reduce wrappers for each supported type and kind _________
   ! _________________________________________________________________________________
 
+    subroutine accelerate(a)
+      real(c_double), intent(in), contiguous :: a(..)
+      call opencoarrays_registernc(c_loc(a),size(a)*c_sizeof(a))  
+    end subroutine
+
     subroutine co_reduce_c_int(a, opr, result_image, stat, errmsg) 
       ! Dummy variables
 #ifdef COMPILER_LACKS_ASSUMED_RANK

src/tests/integration/gpu/co_dot.f90

@@ -11,7 +11,7 @@ module accelerated_module
   ! Explicit interfaces for procedures that wrap accelerated kernels
   interface  
 
-     ! This wrapper exploits the OpenCoarrays acceleration support and is therefore simpler
+     ! This wrapper exploits the OpenCoarrays acceleration support
      subroutine cudaDot(a,b,partial_dot,n) bind(C, name="cudaDot")
        use iso_c_binding, only : c_float,c_int
        real(c_float) :: a(*),b(*)
@@ -80,7 +80,7 @@ program cu_dot_test
   real(c_float) :: dot
   real(c_double) :: t_start, t_end
 
-  ! Library-accelerated variables
+  ! Library-accelerated variables (these are corarrays to facilitate a scatter operation)
   real(c_float), allocatable :: a_acc(:)[:], b_acc(:)[:]
   real(c_float) :: dot_acc[*]
 
@@ -122,10 +122,14 @@ program cu_dot_test
 
   subroutine initialize_all_variables()
     integer(c_int) :: i
-    ! The allocation arguments must be coarrays to support the scatter operation below
-    call accelerated_allocate(a_acc,n_local)
-    call accelerated_allocate(b_acc,n_local)
+    ! These allocation arguments must be coarrays to support the scatter operation below
     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))
+
+    ! Register the desired variables for acceleration
+    call accelerate(a_acc)
+    call accelerate(b_acc)
 
     if(me == 1) then
       ! Initialize the local unaccelerated data on every image