On Fri, Dec 10, 2021 at 05:40:36PM +0000, Kwok Cheung Yeung wrote:
> This adds testcases for metadirectives.
Let me start with the last patch.
> +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-1.c
> @@ -0,0 +1,29 @@
> +/* { dg-do compile } */
> +
> +#define N 100
> +
> +void f (int a[], int b[], int c[])
> +{
> + #pragma omp metadirective \
> + default (teams loop) \
> + default (parallel loop) /* { dg-error "there can only be one default clause in a metadirective before '\\(' token" } */
I'd prefer consistency, check_no_duplicate_clause prints for similar bugs
too many %qs clauses
so it would be nice if this emitted the same (and the before '\\(' token
part would be nice to avoid as well (i.e. use error rather than parse
error).
> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-2.c
> @@ -0,0 +1,74 @@
> +/* { dg-do compile } */
> +
> +#define N 100
> +
> +int main (void)
> +{
> + int x = 0;
> + int y = 0;
> +
> + /* Test implicit default (nothing). */
> + #pragma omp metadirective \
> + when (device={arch("nvptx")}: barrier)
> + x = 1;
I'm not really sure if device={arch("nvptx")} in main is
the best idea for most of such tests.
Because we should be able to decide that right away, main isn't
declare target to (and better shouldn't be) and so when we know host
isn't nvptx and that it won't be offloaded, it is clear it can't be
that arch.
Of course, we need to test such a case too in a few spots, but it would
be nice to have more diversity in the tests.
One possibility is non-main function with declare target to after the
function definition (but one can't then use teams in the metadirectives).
But would be nice to use more variety in selectors, user, implementation, device
with isa or kind, etc. instead of using always the same thing in most of the
tests.
Also it would be nice to cover say a directive which needs loop, a directive
which needs a normal body and say 2 directives which are standalone.
> +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-3.c
> @@ -0,0 +1,31 @@
> +/* { dg-do compile } */
> +/* { dg-additional-options "-fdump-tree-original" } */
> +/* { dg-additional-options "-fdump-tree-gimple" } */
> +/* { dg-additional-options "-fdump-tree-optimized" } */
> +
> +#define N 100
> +
> +void f (int x[], int y[], int z[])
> +{
> + int i;
> +
> + #pragma omp target map(to: x, y) map(from: z)
> + #pragma omp metadirective \
> + when (device={arch("nvptx")}: teams loop) \
> + default (parallel loop)
It would be nice to have many of the tests where all the metadirective
variants are actually possible. Here the nvptx variant
is quite unlikely, nvptx is rarely tested as host arch,
f is not declare target to and even if it was, teams is not allowed
inside of target regions like that.
> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-4.c
> +
> + /* TODO: This does not execute a version of f with the default clause
> + active as might be expected. */
Might be nice to mention that it is correct 5.0 behavior, 5.1 is just
broken in this regard and 5.2 changed the behavior so that parallel loop
is actually invoked.
> + f (a, 2.71828);
> +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-5.c
> @@ -0,0 +1,24 @@
> +/* { dg-do compile } */
> +/* { dg-additional-options "-fdump-tree-original" } */
> +
> +#define N 100
> +
> +void f (int a[], int flag)
> +{
> + int i;
> + #pragma omp metadirective \
> + when (user={condition(flag)}: \
> + target teams distribute parallel for map(from: a[0:N])) \
> + default (parallel for)
> + for (i = 0; i < N; i++)
> + a[i] = i;
> +}
> +
> +/* The metadirective should be resolved at parse time. */
??? How can it? The above is invalid in OpenMP 5.0 (condition
should be constant expression), it is valid in OpenMP 5.1, but is then
resolved at runtime, certainly not at parse time.
Would be nice to also test user={condition(1)} etc. where it would
be resolved at parse time.
And, please add some tests even with user scores.
> +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-6.c
> @@ -0,0 +1,31 @@
> +/* { dg-do compile } */
> +/* { dg-additional-options "-fdump-tree-original" } */
> +/* { dg-additional-options "-fdump-tree-gimple" } */
> +
> +#define N 100
> +
> +void bar (int a[], int run_parallel, int run_guided)
> +{
> + #pragma omp metadirective \
> + when (user={condition(run_parallel)}: parallel)
> + {
> + int i;
> + #pragma omp metadirective \
> + when (construct={parallel}, user={condition(run_guided)}: \
> + for schedule(guided)) \
> + when (construct={parallel}: for schedule(static))
> + for (i = 0; i < N; i++)
> + a[i] = i;
> + }
> + }
> +
> +/* The outer metadirective should be resolved at parse time. */
> +/* The inner metadirective should be resolved during Gimplificiation. */
Again, dynamic condition, so I don't see how this holds, both should be
resolved at runtime.
> +++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-1.c
> @@ -0,0 +1,35 @@
> +/* { dg-do run } */
> +
> +#define N 100
> +
> +void f (int x[], int y[], int z[])
> +{
> + int i;
> +
> + #pragma omp target map(to: x[0:N], y[0:N]) map(from: z[0:N])
> + #pragma omp metadirective \
> + when (device={arch("nvptx")}: teams loop) \
> + default (parallel loop)
> + for (i = 0; i < N; i++)
This doesn't really test which of them was selected and one of them is
extremely unlikely. Doesn't hurt to have some such tests, but would be
nice if there were tests that it could vary (e.g. same test triggering
both or all variants in different code paths, with selector chosen
such that it is possible) where it would return different results
and at runtime it would be possible to decide which one is which.
That can be done either through declare variant in the body, or say
nested metadirective which will say through data sharing or mapping
result in different values (say task shared(var) in one case and
task firstprivate(var) in another etc.).
Jakub
new file mode 100644
@@ -0,0 +1,29 @@
+/* { dg-do compile } */
+
+#define N 100
+
+void f (int a[], int b[], int c[])
+{
+ #pragma omp metadirective \
+ default (teams loop) \
+ default (parallel loop) /* { dg-error "there can only be one default clause in a metadirective before '\\(' token" } */
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+ #pragma omp metadirective \
+ default (bad_directive) /* { dg-error "unknown directive name before '\\)' token" } */
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+ #pragma omp metadirective \
+ default (teams loop) \
+ where (device={arch("nvptx")}: parallel loop) /* { dg-error "expected 'when' or 'default' before '\\(' token" } */
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+ #pragma omp metadirective \
+ default (teams loop) \
+ when (device={arch("nvptx")} parallel loop) /* { dg-error "expected colon before 'parallel'" } */
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+ #pragma omp metadirective \
+ default (metadirective default (flush)) /* { dg-error "metadirectives cannot be used as directive variants before 'default'" } */
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+}
new file mode 100644
@@ -0,0 +1,74 @@
+/* { dg-do compile } */
+
+#define N 100
+
+int main (void)
+{
+ int x = 0;
+ int y = 0;
+
+ /* Test implicit default (nothing). */
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: barrier)
+ x = 1;
+
+ /* Test with multiple standalone directives. */
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: barrier) \
+ default (flush)
+ x = 1;
+
+ /* Test combining a standalone directive with one that takes a statement
+ body. */
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: parallel) \
+ default (barrier)
+ x = 1;
+
+ /* Test combining a standalone directive with one that takes a for loop. */
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: parallel for) \
+ default (barrier)
+ for (int i = 0; i < N; i++)
+ x += i;
+
+ /* Test combining a directive that takes a for loop with one that takes
+ a regular statement body. */
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: parallel for) \
+ default (parallel)
+ for (int i = 0; i < N; i++)
+ x += i;
+
+ /* Test labels inside statement body. */
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: teams num_teams(512)) \
+ when (device={arch("gcn")}: teams num_teams(256)) \
+ default (teams num_teams(4))
+ {
+ if (x)
+ goto l1;
+ else
+ goto l2;
+ l1: ;
+ l2: ;
+ }
+
+ /* Test local labels inside statement body. */
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: teams num_teams(512)) \
+ when (device={arch("gcn")}: teams num_teams(256)) \
+ default (teams num_teams(4))
+ {
+ //__label__ l1, l2;
+
+ if (x)
+ goto l1;
+ else
+ goto l2;
+ l1: ;
+ l2: ;
+ }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,31 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+/* { dg-additional-options "-fdump-tree-optimized" } */
+
+#define N 100
+
+void f (int x[], int y[], int z[])
+{
+ int i;
+
+ #pragma omp target map(to: x, y) map(from: z)
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: teams loop) \
+ default (parallel loop)
+ for (i = 0; i < N; i++)
+ z[i] = x[i] * y[i];
+}
+
+/* The metadirective should be resolved after Gimplification. */
+
+/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "when \\(device arch .nvptx.\\):" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "default:" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp loop" 2 "original" } } */
+
+/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "gimple" } } */
+
+/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "optimized" } } */
new file mode 100644
@@ -0,0 +1,40 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 100
+
+#pragma omp declare target
+void f(double a[], double x) {
+ int i;
+
+ #pragma omp metadirective \
+ when (construct={target}: distribute parallel for) \
+ default (parallel for simd)
+ for (i = 0; i < N; i++)
+ a[i] = x * i;
+}
+#pragma omp end declare target
+
+ int main()
+{
+ double a[N];
+
+ #pragma omp target teams map(from: a[0:N])
+ f (a, 3.14159);
+
+ /* TODO: This does not execute a version of f with the default clause
+ active as might be expected. */
+ f (a, 2.71828);
+
+ return 0;
+ }
+
+ /* The metadirective should be resolved during Gimplification. */
+
+/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "when \\(construct target.*\\):" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "default:" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } } */
+
+/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */
new file mode 100644
@@ -0,0 +1,24 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+#define N 100
+
+void f (int a[], int flag)
+{
+ int i;
+ #pragma omp metadirective \
+ when (user={condition(flag)}: \
+ target teams distribute parallel for map(from: a[0:N])) \
+ default (parallel for)
+ for (i = 0; i < N; i++)
+ a[i] = i;
+}
+
+/* The metadirective should be resolved at parse time. */
+
+/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp distribute" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp for" 2 "original" } } */
new file mode 100644
@@ -0,0 +1,31 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 100
+
+void bar (int a[], int run_parallel, int run_guided)
+{
+ #pragma omp metadirective \
+ when (user={condition(run_parallel)}: parallel)
+ {
+ int i;
+ #pragma omp metadirective \
+ when (construct={parallel}, user={condition(run_guided)}: \
+ for schedule(guided)) \
+ when (construct={parallel}: for schedule(static))
+ for (i = 0; i < N; i++)
+ a[i] = i;
+ }
+ }
+
+/* The outer metadirective should be resolved at parse time. */
+/* The inner metadirective should be resolved during Gimplificiation. */
+
+/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 2 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp for" 4 "original" } } */
+/* { dg-final { scan-tree-dump-times "when \\(construct parallel" 4 "original" } } */
+/* { dg-final { scan-tree-dump-times "default:" 2 "original" } } */
+
+/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */
new file mode 100644
@@ -0,0 +1,15 @@
+int main (void)
+{
+ int x, y;
+
+ /* Test nested functions inside statement body. */
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: teams num_teams(512)) \
+ when (device={arch("gcn")}: teams num_teams(256)) \
+ default (teams num_teams(4))
+ {
+ int f (int x) { return x * 3; }
+
+ y = f (x);
+ }
+}
new file mode 100644
@@ -0,0 +1,41 @@
+! { dg-do compile }
+
+program main
+ integer, parameter :: N = 10
+ integer, dimension(N) :: a
+ integer, dimension(N) :: b
+ integer, dimension(N) :: c
+ integer :: i
+
+ do i = 1, N
+ a(i) = i * 2
+ b(i) = i * 3
+ end do
+
+ !$omp metadirective &
+ !$omp& default (teams loop) &
+ !$omp& default (parallel loop) ! { dg-error "there can only be one default clause in a metadirective at .1." }
+ do i = 1, N
+ c(i) = a(i) * b(i)
+ end do
+
+ !$omp metadirective default (xyz) ! { dg-error "Unclassifiable OpenMP directive at .1." }
+ do i = 1, N
+ c(i) = a(i) * b(i)
+ end do
+
+ !$omp metadirective &
+ !$omp& default (teams loop) & ! { dg-error "expected 'default' or 'when' at .1." }
+ !$omp& where (device={arch("nvptx")}: parallel loop)
+ do i = 1, N
+ c(i) = a(i) * b(i)
+ end do
+
+ !$omp begin metadirective &
+ !$omp& when (device={arch("nvptx")}: parallel do) &
+ !$omp& default (barrier) ! { dg-error "variant directive used in OMP BEGIN METADIRECTIVE at .1. must have a corresponding end directive" }
+ do i = 1, N
+ c(i) = a(i) * b(i)
+ end do
+ !$omp end metadirective ! { dg-error "Unexpected !OMP END METADIRECTIVE statement at .1." }
+end program
new file mode 100644
@@ -0,0 +1,59 @@
+! { dg-do compile }
+
+program main
+ integer, parameter :: N = 100
+ integer :: x = 0
+ integer :: y = 0
+ integer :: i
+
+ ! Test implicit default directive
+ !$omp metadirective &
+ !$omp& when (device={arch("nvptx")}: barrier)
+ x = 1
+
+ ! Test implicit default directive combined with a directive that takes a
+ ! do loop.
+ !$omp metadirective &
+ !$omp& when (device={arch("nvptx")}: parallel do)
+ do i = 1, N
+ x = x + i
+ end do
+
+ ! Test with multiple standalone directives.
+ !$omp metadirective &
+ !$omp& when (device={arch("nvptx")}: barrier) &
+ !$omp& default (flush)
+ x = 1
+
+ ! Test combining a standalone directive with one that takes a do loop.
+ !$omp metadirective &
+ !$omp& when (device={arch("nvptx")}: parallel do) &
+ !$omp& default (barrier)
+ do i = 1, N
+ x = x + i
+ end do
+
+ ! Test combining a directive that takes a do loop with one that takes
+ ! a statement body.
+ !$omp begin metadirective &
+ !$omp& when (device={arch("nvptx")}: parallel do) &
+ !$omp& default (parallel)
+ do i = 1, N
+ x = x + i
+ end do
+ !$omp end metadirective
+
+ ! Test labels in the body
+ !$omp begin metadirective &
+ !$omp& when (device={arch("nvptx")}: parallel do) &
+ !$omp& when (device={arch("gcn")}: parallel)
+ do i = 1, N
+ x = x + i
+ if (x .gt. N/2) goto 10
+10 x = x + 1
+ goto 20
+ x = x + 2
+20 continue
+ end do
+ !$omp end metadirective
+end program
new file mode 100644
@@ -0,0 +1,34 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-additional-options "-fdump-tree-gimple" }
+! { dg-additional-options "-fdump-tree-optimized" }
+
+module test
+ integer, parameter :: N = 100
+contains
+ subroutine f (x, y, z)
+ integer :: x(N), y(N), z(N)
+
+ !$omp target map (to: v1, v2) map(from: v3)
+ !$omp metadirective &
+ !$omp& when(device={arch("nvptx")}: teams loop) &
+ !$omp& default(parallel loop)
+ do i = 1, N
+ z(i) = x(i) * y(i)
+ enddo
+ !$omp end target
+ end subroutine
+end module
+
+! The metadirective should be resolved after Gimplification.
+
+! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } }
+! { dg-final { scan-tree-dump-times "when \\(device arch .nvptx.\\):" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "original" } }
+! { dg-final { scan-tree-dump-times "default:" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp loop" 2 "original" } }
+
+! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "gimple" } }
+
+! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "optimized" } }
new file mode 100644
@@ -0,0 +1,39 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program test
+ implicit none
+ integer, parameter :: N = 100
+ real :: a(N)
+
+ !$omp target map(from: a)
+ call f (a, 3.14159)
+ !$omp end target
+
+ ! TODO: This does not execute a version of f with the default clause
+ ! active as might be expected.
+ call f (a, 2.71828)
+contains
+ subroutine f (a, x)
+ integer :: i
+ real :: a(N), x
+ !$omp declare target
+
+ !$omp metadirective &
+ !$omp& when (construct={target}: distribute parallel do ) &
+ !$omp& default(parallel do simd)
+ do i = 1, N
+ a(i) = x * i
+ end do
+ end subroutine
+end program
+
+! The metadirective should be resolved during Gimplification.
+
+! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } }
+! { dg-final { scan-tree-dump-times "when \\(construct target.*\\):" 1 "original" } }
+! { dg-final { scan-tree-dump-times "default:" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } }
+
+! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } }
new file mode 100644
@@ -0,0 +1,30 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module test
+ integer, parameter :: N = 100
+contains
+ subroutine f (a, flag)
+ integer :: a(N)
+ logical :: flag
+ integer :: i
+
+ !$omp metadirective &
+ !$omp& when (user={condition(flag)}: &
+ !$omp& target teams distribute parallel do map(from: a(1:N))) &
+ !$omp& default(parallel do)
+ do i = 1, N
+ a(i) = i
+ end do
+ end subroutine
+end module
+
+! The metadirective should be resolved at parse time, but is currently
+! resolved during Gimplification
+
+! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp distribute" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp for" 2 "gimple" } }
new file mode 100644
@@ -0,0 +1,31 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module test
+ integer, parameter :: N = 100
+contains
+ subroutine f (a, run_parallel, run_guided)
+ integer :: a(N)
+ logical :: run_parallel, run_guided
+ integer :: i
+
+ !$omp begin metadirective when(user={condition(run_parallel)}: parallel)
+ !$omp metadirective &
+ !$omp& when(construct={parallel}, user={condition(run_guided)}: &
+ !$omp& do schedule(guided)) &
+ !$omp& when(construct={parallel}: do schedule(static))
+ do i = 1, N
+ a(i) = i
+ end do
+ !$omp end metadirective
+ end subroutine
+end module
+
+! The outer metadirective should be resolved at parse time, but is
+! currently resolved during Gimplification.
+
+! The inner metadirective should be resolved during Gimplificiation.
+
+! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp for" 2 "gimple" } }
new file mode 100644
@@ -0,0 +1,35 @@
+/* { dg-do run } */
+
+#define N 100
+
+void f (int x[], int y[], int z[])
+{
+ int i;
+
+ #pragma omp target map(to: x[0:N], y[0:N]) map(from: z[0:N])
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: teams loop) \
+ default (parallel loop)
+ for (i = 0; i < N; i++)
+ z[i] = x[i] * y[i];
+}
+
+int main (void)
+{
+ int x[N], y[N], z[N];
+ int i;
+
+ for (i = 0; i < N; i++)
+ {
+ x[i] = i;
+ y[i] = -i;
+ }
+
+ f (x, y, z);
+
+ for (i = 0; i < N; i++)
+ if (z[i] != x[i] * y[i])
+ return 1;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,41 @@
+/* { dg-do run } */
+
+#include <math.h>
+
+#define N 100
+#define EPSILON 0.001
+
+#pragma omp declare target
+void f(double a[], double x) {
+ int i;
+
+ #pragma omp metadirective \
+ when (construct={target}: distribute parallel for) \
+ default (parallel for simd)
+ for (i = 0; i < N; i++)
+ a[i] = x * i;
+}
+#pragma omp end declare target
+
+ int main()
+{
+ double a[N];
+ int i;
+
+ #pragma omp target teams map(from: a[0:N])
+ f (a, M_PI);
+
+ for (i = 0; i < N; i++)
+ if (fabs (a[i] - (M_PI * i)) > EPSILON)
+ return 1;
+
+ /* TODO: This does not execute a version of f with the default clause
+ active as might be expected. */
+ f (a, M_E);
+
+ for (i = 0; i < N; i++)
+ if (fabs (a[i] - (M_E * i)) > EPSILON)
+ return 1;
+
+ return 0;
+ }
new file mode 100644
@@ -0,0 +1,34 @@
+/* { dg-do run } */
+
+#define N 100
+
+int f (int a[], int flag)
+{
+ int i;
+ int res = 0;
+
+ #pragma omp metadirective \
+ when (user={condition(!flag)}: \
+ target teams distribute parallel for \
+ map(from: a[0:N]) private(res)) \
+ default (parallel for)
+ for (i = 0; i < N; i++)
+ {
+ a[i] = i;
+ res = 1;
+ }
+
+ return res;
+}
+
+int main (void)
+{
+ int a[N];
+
+ if (f (a, 0))
+ return 1;
+ if (!f (a, 1))
+ return 1;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,52 @@
+/* { dg-do run } */
+
+#include <omp.h>
+
+#define N 100
+
+int f (int a[], int run_parallel, int run_static)
+{
+ int is_parallel = 0;
+ int is_static = 0;
+
+ #pragma omp metadirective \
+ when (user={condition(run_parallel)}: parallel)
+ {
+ int i;
+
+ if (omp_in_parallel ())
+ is_parallel = 1;
+
+ #pragma omp metadirective \
+ when (construct={parallel}, user={condition(!run_static)}: \
+ for schedule(guided) private(is_static)) \
+ when (construct={parallel}: for schedule(static))
+ for (i = 0; i < N; i++)
+ {
+ a[i] = i;
+ is_static = 1;
+ }
+ }
+
+ return (is_parallel << 1) | is_static;
+}
+
+int main (void)
+{
+ int a[N];
+
+ /* is_static is always set if run_parallel is false. */
+ if (f (a, 0, 0) != 1)
+ return 1;
+
+ if (f (a, 0, 1) != 1)
+ return 1;
+
+ if (f (a, 1, 0) != 2)
+ return 1;
+
+ if (f (a, 1, 1) != 3)
+ return 1;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,33 @@
+! { dg-do run }
+
+program test
+ implicit none
+
+ integer, parameter :: N = 100
+ integer :: x(N), y(N), z(N)
+ integer :: i
+
+ do i = 1, N
+ x(i) = i;
+ y(i) = -i;
+ end do
+
+ call f (x, y, z)
+
+ do i = 1, N
+ if (z(i) .ne. x(i) * y(i)) stop 1
+ end do
+contains
+ subroutine f (x, y, z)
+ integer :: x(N), y(N), z(N)
+
+ !$omp target map (to: x, y) map(from: z)
+ !$omp metadirective &
+ !$omp& when(device={arch("nvptx")}: teams loop) &
+ !$omp& default(parallel loop)
+ do i = 1, N
+ z(i) = x(i) * y(i)
+ enddo
+ !$omp end target
+ end subroutine
+end program
new file mode 100644
@@ -0,0 +1,40 @@
+! { dg-do run }
+
+program test
+ implicit none
+ integer, parameter :: N = 100
+ real, parameter :: PI_CONST = 3.14159
+ real, parameter :: E_CONST = 2.71828
+ real, parameter :: EPSILON = 0.001
+ integer :: i
+ real :: a(N)
+
+ !$omp target map(from: a)
+ call f (a, PI_CONST)
+ !$omp end target
+
+ do i = 1, N
+ if (abs (a(i) - (PI_CONST * i)) .gt. EPSILON) stop 1
+ end do
+
+ ! TODO: This does not execute a version of f with the default clause
+ ! active as might be expected.
+ call f (a, E_CONST)
+
+ do i = 1, N
+ if (abs (a(i) - (E_CONST * i)) .gt. EPSILON) stop 2
+ end do
+contains
+ subroutine f (a, x)
+ integer :: i
+ real :: a(N), x
+ !$omp declare target
+
+ !$omp metadirective &
+ !$omp& when (construct={target}: distribute parallel do ) &
+ !$omp& default(parallel do simd)
+ do i = 1, N
+ a(i) = x * i
+ end do
+ end subroutine
+end program
new file mode 100644
@@ -0,0 +1,29 @@
+! { dg-do run }
+
+program test
+ implicit none
+
+ integer, parameter :: N = 100
+ integer :: a(N)
+ integer :: res
+
+ if (f (a, .false.)) stop 1
+ if (.not. f (a, .true.)) stop 2
+contains
+ logical function f (a, flag)
+ integer :: a(N)
+ logical :: flag
+ logical :: res = .false.
+ integer :: i
+ f = .false.
+ !$omp metadirective &
+ !$omp& when (user={condition(.not. flag)}: &
+ !$omp& target teams distribute parallel do &
+ !$omp& map(from: a(1:N)) private(res)) &
+ !$omp& default(parallel do)
+ do i = 1, N
+ a(i) = i
+ f = .true.
+ end do
+ end function
+end program
new file mode 100644
@@ -0,0 +1,46 @@
+! { dg-do run }
+
+program test
+ use omp_lib
+
+ implicit none
+ integer, parameter :: N = 100
+ integer :: a(N)
+ logical :: is_parallel, is_static
+
+ ! is_static is always set if run_parallel is false.
+ call f (a, .false., .false., is_parallel, is_static)
+ if (is_parallel .or. .not. is_static) stop 1
+
+ call f (a, .false., .true., is_parallel, is_static)
+ if (is_parallel .or. .not. is_static) stop 2
+
+ call f (a, .true., .false., is_parallel, is_static)
+ if (.not. is_parallel .or. is_static) stop 3
+
+ call f (a, .true., .true., is_parallel, is_static)
+ if (.not. is_parallel .or. .not. is_static) stop 4
+contains
+ subroutine f (a, run_parallel, run_static, is_parallel, is_static)
+ integer :: a(N)
+ logical, intent(in) :: run_parallel, run_static
+ logical, intent(out) :: is_parallel, is_static
+ integer :: i
+
+ is_parallel = .false.
+ is_static = .false.
+
+ !$omp begin metadirective when(user={condition(run_parallel)}: parallel)
+ if (omp_in_parallel ()) is_parallel = .true.
+
+ !$omp metadirective &
+ !$omp& when(construct={parallel}, user={condition(.not. run_static)}: &
+ !$omp& do schedule(guided) private(is_static)) &
+ !$omp& when(construct={parallel}: do schedule(static))
+ do i = 1, N
+ a(i) = i
+ is_static = .true.
+ end do
+ !$omp end metadirective
+ end subroutine
+end program