Hi, The attached patch adds additional test for the routine directive for C/C++/Fortran.
Committed to gomp-4_0-branch. Jim
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c new file mode 100644 index 0000000..73ca528 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c @@ -0,0 +1,32 @@ + +/* { dg-do run } */ + +#include <stdlib.h> + +#pragma acc routine nohost +int +foo (int n) +{ + if (n == 0 || n == 1) + return 1; + + return n * n; +} + +int +main() +{ + int a, n = 10; + +#pragma acc parallel copy (a, n) + { + a = foo (n); + } + + if (a != n * n) + abort (); + + return 0; +} + +/* { dg-output "foo not found" { target openacc_host_selected } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-4.c new file mode 100644 index 0000000..c73400c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-4.c @@ -0,0 +1,127 @@ + +/* { dg-do run } */ + +#include <stdlib.h> + +#define M 8 +#define N 32 + +#pragma acc routine vector +void +vector (int *a) +{ + int i; + +#pragma acc loop vector + for (i = 0; i < N; i++) + a[i] -= a[i]; +} + +#pragma acc routine worker +void +worker (int *b) +{ + int i, j; + +#pragma acc loop gang + for (i = 0; i < N; i++) + { +#pragma acc loop worker + for (j = 0; j < M; j++) + b[i * M + j] += b[i * M + j]; + } +} + +#pragma acc routine gang +void +gang (int *a) +{ + int i; + +#pragma acc loop gang + for (i = 0; i < N; i++) + a[i] -= i; +} + +#pragma acc routine seq +void +seq (int *a) +{ + int i; + + for (i = 0; i < N; i++) + a[i] += 1; +} + +#include <stdio.h> + +int +main(int argc, char **argv) +{ + int i; + int a[N]; + int b[M * N]; + + i = 0; + + for (i = 0; i < N; i++) + a[i] = 0; + +#pragma acc parallel copy (a[0:N]) + { +#pragma acc loop seq + for (i = 0; i < N; i++) + seq (&a[0]); + } + + for (i = 0; i < N; i++) + { + if (a[i] != N) + abort (); + } + +#pragma acc parallel copy (a[0:N]) + { +#pragma acc loop seq + for (i = 0; i < N; i++) + gang (&a[0]); + } + + for (i = 0; i < N; i++) + { + if (a[i] != N + (N * (-1 * i))) + abort (); + } + + for (i = 0; i < N; i++) + a[i] = i; + +#pragma acc parallel copy (b[0:M*N]) + { + worker (&b[0]); + } + + for (i = 0; i < N; i++) + { + if (a[i] != i) + abort (); + } + + for (i = 0; i < N; i++) + a[i] = i; + +#pragma acc parallel copy (a[0:N]) + { +#pragma acc loop vector + for (i = 0; i < N; i++) + vector (&a[0]); + } + + for (i = 0; i < N; i++) + { + if (a[i] != 0) + abort (); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-5.c new file mode 100644 index 0000000..6d0fbe3 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-5.c @@ -0,0 +1,62 @@ + +/* { dg-do run } */ + +#include <stdio.h> +#include <stdlib.h> + +#pragma acc routine bind (foo) +int +subr1 (int n) +{ + if (n == 0 || n == 1) + return 1; + + return n * foo (n - 1); +} + +#pragma acc routine bind ("bar") +int +subr2 (int n) +{ + if (n == 0 || n == 1) + return 1; + + return n * bar (n - 1); +} + +int +main() +{ + int *a, i, n = 10; + + a = (int *)malloc (sizeof (int) * n); + +#pragma acc parallel copy (a[0:n]) vector_length (5) + { +#pragma acc loop + for (i = 0; i < n; i++) + a[i] = foo (i); + } + + for (i = 0; i < n; i++) + if (a[i] != subr1 (i)) + abort (); + + for (i = 0; i < n; i++) + a[i] = 0; + +#pragma acc parallel copy (a[0:n]) vector_length (5) + { +#pragma acc loop + for (i = 0; i < n; i++) + a[i] = bar (i); + } + + for (i = 0; i < n; i++) + if (a[i] != subr2 (i)) + abort (); + + free (a); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-5.f90 index aaeb994..956da8e 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/routine-5.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-5.f90 @@ -15,7 +15,7 @@ program main contains function func (n) result (rc) - !$acc routine gang worker vector seq nohost + !$acc routine integer, intent (in) :: n integer :: rc diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90 new file mode 100644 index 0000000..4b7b707 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90 @@ -0,0 +1,28 @@ +! { dg-do run } + +program main + integer :: a, n + + n = 10 + + !$acc parallel copy (a, n) + a = foo (n) + !$acc end parallel + + if (a .ne. n * n) call abort + +contains + +function foo (n) result (rc) + !$acc routine nohost + + integer, intent (in) :: n + integer :: rc + + rc = n * n + +end function + +end program main + +! { dg-output "not found" { target openacc_host_selected } } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 new file mode 100644 index 0000000..354784e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 @@ -0,0 +1,121 @@ + +! { dg-do run } +! { dg-additional-options "-cpp" } + +#define M 8 +#define N 32 + +program main + integer :: i + integer :: a(N) + integer :: b(M * N) + + do i = 1, N + a(i) = 0 + end do + + !$acc parallel copy (a) + !$acc loop seq + do i = 1, N + call seq (a) + end do + !$acc end parallel + + do i = 1, N + if (a(i) .ne.N) call abort + end do + + !$acc parallel copy (a) + !$acc loop seq + do i = 1, N + call gang (a) + end do + !$acc end parallel + + do i = 1, N + if (a(i) .ne. (N + (N * (-1 * i)))) call abort + end do + + do i = 1, N + b(i) = i + end do + + !$acc parallel copy (b) + !$acc loop worker + do i = 1, N + call worker (b) + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. N + i) call abort + end do + + do i = 1, N + a(i) = i + end do + + !$acc parallel copy (a) + !$acc loop vector + do i = 1, N + call vector (a) + end do + !$acc end parallel + + do i = 1, N + if (a(i) .ne. 0) call abort + end do + +contains + +subroutine vector (a) + !$acc routine vector + integer, intent (inout) :: a(N) + integer :: i + + !$acc loop vector + do i = 1, N + a(i) = a(i) - a(i) + end do + +end subroutine vector + +subroutine worker (b) + !$acc routine worker + integer, intent (inout) :: b(M*N) + integer :: i, j + + !$acc loop gang + do i = 1, N + !$acc loop worker + do j = 1, M + b(j + ((i - 1) * M)) = b(j + ((i - 1) * M)) + 1 + end do + end do + +end subroutine worker + +subroutine gang (a) + !$acc routine gang + integer, intent (inout) :: a(N) + integer :: i + + !$acc loop gang + do i = 1, N + a(i) = a(i) - i + end do + +end subroutine gang + +subroutine seq (a) + !$acc routine seq + integer, intent (inout) :: a(M) + integer :: i + + do i = 1, N + a(i) = a(i) + 1 + end do + +end subroutine seq + +end program main diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-8.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-8.f90 new file mode 100644 index 0000000..2060740 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-8.f90 @@ -0,0 +1,61 @@ + +! { dg-do run } + +program main + integer, parameter :: n = 10 + integer :: a(n) + integer :: i + + !$acc parallel copy (a) vector_length (5) + !$acc loop + do i = 1, n + a(i) = foo (i); + end do + !$acc end parallel + + do i = 1, n + if (a(i) .ne. subr1 (i)) call abort + end do + + do i = 1, n + a(i) = 0 + end do + + !$acc parallel copy (a) vector_length (5) + !$acc loop + do i = 1, n + a(i) = bar (i); + end do + !$acc end parallel + + do i = 1, n + if (a(i) .ne. subr2 (i)) call abort + end do + +contains + +function subr1 (n) result (rc) + !$acc routine bind (foo) + integer :: n, rc + + if ((n .eq. 0) .or. (n .eq. 1)) then + rc = 1 + else + rc = n * foo (n - 1); + end if + +end function + +function subr2 (n) result (rc) + !$acc routine bind ("bar") + integer :: n, rc + + if ((n .eq. 0) .or. (n .eq. 1)) then + rc = 1 + else + rc = n * bar (n - 1); + end if + +end function + +end program main