31dd69b7ff
gcc/testsuite/ * c-c++-common/goacc/deviceptr-4.c: New file. * c-c++-common/goacc/kernels-counter-var-redundant-load.c: Likewise. * c-c++-common/goacc/kernels-loop-data-2.c: Likewise. * c-c++-common/goacc/kernels-loop-data-enter-exit-2.c: Likewise. * c-c++-common/goacc/kernels-loop-data-enter-exit.c: Likewise. * c-c++-common/goacc/kernels-loop-data-update.c: Likewise. * c-c++-common/goacc/kernels-loop-data.c: Likewise. * c-c++-common/goacc/kernels-parallel-loop-data-enter-exit.c: Likewise. * c-c++-common/goacc/parallel-reduction.c: Likewise. * c-c++-common/goacc/private-reduction-1.c: Likewise. * gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95: Likewise. * gfortran.dg/goacc/modules.f95: Likewise. * gfortran.dg/goacc/routine-8.f90: Likewise. * gfortran.dg/goacc/routine-level-of-parallelism-1.f90: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c: Don't force "-O2". * testsuite/libgomp.oacc-c-c++-common/data-2.c: Update. * testsuite/libgomp.oacc-c-c++-common/host_data-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/mode-transitions.c: Likewise. * testsuite/libgomp.oacc-fortran/data-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/data-2.f90: Likewise. * testsuite/libgomp.oacc-c++/non-scalar-data.C: New file. * testsuite/libgomp.oacc-c-c++-common/declare-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/enter-data.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-data-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-data-enter-exit-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-data-enter-exit.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-data-update.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-data.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-parallel-loop-data-enter-exit.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-6.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-6.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-7.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-reduction-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-loop-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-loop-1.h: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-loop-2.h: Likewise. * testsuite/libgomp.oacc-fortran/cublas-fixed.h: Likewise. * testsuite/libgomp.oacc-fortran/dummy-array.f90: Likewise. * testsuite/libgomp.oacc-fortran/host_data-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/host_data-3.f: Likewise. * testsuite/libgomp.oacc-fortran/host_data-4.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-acc-loop-reduction-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-acc-loop-reduction.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-collapse-3.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-collapse-4.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-independent.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-loop-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-map-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-parallel-loop-data-enter-exit.f95: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-3.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-6.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-3.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-4.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-5.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-6.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-7.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-reduction-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/lib-12.f90: Likewise. * testsuite/libgomp.oacc-fortran/lib-13.f90: Likewise. * testsuite/libgomp.oacc-fortran/lib-14.f90: Likewise. * testsuite/libgomp.oacc-fortran/lib-15.f90: Likewise. * testsuite/libgomp.oacc-fortran/parallel-loop-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/reference-reductions.f90: Likewise. * testsuite/libgomp.oacc-fortran/vector-routine.f90: Likewise. Co-Authored-By: James Norris <jnorris@codesourcery.com> Co-Authored-By: Julian Brown <julian@codesourcery.com> Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com> Co-Authored-By: Tom de Vries <tom@codesourcery.com> From-SVN: r261884
99 lines
2.1 KiB
Fortran
99 lines
2.1 KiB
Fortran
! Test host_data interoperability with CUDA blas. This test was
|
|
! derived from libgomp.oacc-c-c++-common/host_data-1.c.
|
|
|
|
! { dg-do run { target openacc_nvidia_accel_selected } }
|
|
! { dg-additional-options "-lcublas -Wall -Wextra" }
|
|
|
|
program test
|
|
implicit none
|
|
|
|
integer, parameter :: N = 10
|
|
integer :: i
|
|
real*4 :: x_ref(N), y_ref(N), x(N), y(N), a
|
|
|
|
interface
|
|
subroutine cublassaxpy(N, alpha, x, incx, y, incy) bind(c, name="cublasSaxpy")
|
|
use iso_c_binding
|
|
integer(kind=c_int), value :: N
|
|
real(kind=c_float), value :: alpha
|
|
type(*), dimension(*) :: x
|
|
integer(kind=c_int), value :: incx
|
|
type(*), dimension(*) :: y
|
|
integer(kind=c_int), value :: incy
|
|
end subroutine cublassaxpy
|
|
end interface
|
|
|
|
a = 2.0
|
|
|
|
do i = 1, N
|
|
x(i) = 4.0 * i
|
|
y(i) = 3.0
|
|
x_ref(i) = x(i)
|
|
y_ref(i) = y(i)
|
|
end do
|
|
|
|
call saxpy (N, a, x_ref, y_ref)
|
|
|
|
!$acc data copyin (x) copy (y)
|
|
!$acc host_data use_device (x, y)
|
|
call cublassaxpy(N, a, x, 1, y, 1)
|
|
!$acc end host_data
|
|
!$acc end data
|
|
|
|
call validate_results (N, y, y_ref)
|
|
|
|
!$acc data create (x) copyout (y)
|
|
!$acc parallel loop
|
|
do i = 1, N
|
|
y(i) = 3.0
|
|
end do
|
|
!$acc end parallel loop
|
|
|
|
!$acc host_data use_device (x, y)
|
|
call cublassaxpy(N, a, x, 1, y, 1)
|
|
!$acc end host_data
|
|
!$acc end data
|
|
|
|
call validate_results (N, y, y_ref)
|
|
|
|
y(:) = 3.0
|
|
|
|
!$acc data copyin (x) copyin (a) copy (y)
|
|
!$acc parallel present (x) pcopy (y) present (a)
|
|
call saxpy (N, a, x, y)
|
|
!$acc end parallel
|
|
!$acc end data
|
|
|
|
call validate_results (N, y, y_ref)
|
|
|
|
y(:) = 3.0
|
|
|
|
!$acc enter data copyin (x, a, y)
|
|
!$acc parallel present (x) pcopy (y) present (a)
|
|
call saxpy (N, a, x, y)
|
|
!$acc end parallel
|
|
!$acc exit data delete (x, a) copyout (y)
|
|
|
|
call validate_results (N, y, y_ref)
|
|
end program test
|
|
|
|
subroutine saxpy (nn, aa, xx, yy)
|
|
integer :: nn
|
|
real*4 :: aa, xx(nn), yy(nn)
|
|
integer i
|
|
!$acc routine
|
|
|
|
do i = 1, nn
|
|
yy(i) = yy(i) + aa * xx(i)
|
|
end do
|
|
end subroutine saxpy
|
|
|
|
subroutine validate_results (n, a, b)
|
|
integer :: n
|
|
real*4 :: a(n), b(n)
|
|
|
|
do i = 1, N
|
|
if (abs(a(i) - b(i)) > 0.0001) call abort
|
|
end do
|
|
end subroutine validate_results
|