procyberian/gcc
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

OpenMP: Shared metadirective/dynamic selector tests for C and C++

Browse source 
gcc/testsuite/ChangeLog
	* c-c++-common/gomp/adjust-args-6.c: New.
	* c-c++-common/gomp/attrs-metadirective-1.c: New.
	* c-c++-common/gomp/attrs-metadirective-2.c: New.
	* c-c++-common/gomp/attrs-metadirective-3.c: New.
	* c-c++-common/gomp/attrs-metadirective-4.c: New.
	* c-c++-common/gomp/attrs-metadirective-5.c: New.
	* c-c++-common/gomp/attrs-metadirective-6.c: New.
	* c-c++-common/gomp/attrs-metadirective-7.c: New.
	* c-c++-common/gomp/attrs-metadirective-8.c: New.
	* c-c++-common/gomp/declare-variant-arg-exprs.c: New.
	* c-c++-common/gomp/declare-variant-dynamic-1.c: New.
	* c-c++-common/gomp/declare-variant-dynamic-2.c: New.
	* c-c++-common/gomp/metadirective-1.c: New.
	* c-c++-common/gomp/metadirective-2.c: New.
	* c-c++-common/gomp/metadirective-3.c: New.
	* c-c++-common/gomp/metadirective-4.c: New.
	* c-c++-common/gomp/metadirective-5.c: New.
	* c-c++-common/gomp/metadirective-6.c: New.
	* c-c++-common/gomp/metadirective-7.c: New.
	* c-c++-common/gomp/metadirective-8.c: New.
	* c-c++-common/gomp/metadirective-construct.c: New.
	* c-c++-common/gomp/metadirective-device.c: New.
	* c-c++-common/gomp/metadirective-no-score.c: New.
	* c-c++-common/gomp/metadirective-target-device-1.c: New.
	* c-c++-common/gomp/metadirective-target-device-2.c: New.

libgomp/ChangeLog
	* testsuite/libgomp.c-c++-common/metadirective-1.c: New.
	* testsuite/libgomp.c-c++-common/metadirective-2.c: New.
	* testsuite/libgomp.c-c++-common/metadirective-3.c: New.
	* testsuite/libgomp.c-c++-common/metadirective-4.c: New.
	* testsuite/libgomp.c-c++-common/metadirective-5.c: New.
	* testsuite/libgomp.c-c++-common/metadirective-late-1.c: New.
	* testsuite/libgomp.c-c++-common/metadirective-late-2.c: New.
	* testsuite/libgomp.c-c++-common/metadirective-target-device.c: New.

Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>
Co-Authored-By: Sandra Loosemore <sandra@codesourcery.com>
This commit is contained in:
Sandra Loosemore 2025-01-08 01:55:47 +00:00
parent 677e452e55
commit fdeceba59b
33 changed files with 1858 additions and 0 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

26
gcc/testsuite/c-c++-common/gomp/adjust-args-6.c Normal file
View file

@ -0,0 +1,26 @@
/* { dg-do compile } */
/* { dg-additional-options "-fdump-tree-gimple" } */
/* Ensure that adjust_args is only applied when variant substitution happens. */
extern int flag;
void h(int *);
void f(int *);
#pragma omp declare variant(f) match(construct={dispatch}, user={condition(flag)}) adjust_args(need_device_ptr : x)
void g(int *x);
void foo(int *y)
{
#pragma omp dispatch
h(y);
#pragma omp dispatch
f(y);
#pragma omp dispatch
g(y); /* { dg-bogus "late or dynamic variant resolution" "" { xfail *-*-* } } */
}
/* { dg-final { scan-tree-dump-times "h \\(y\\);" 1 "gimple" } } */
/* { dg-final { scan-tree-dump-times "f \\(y\\);" 2 "gimple" } } */
/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(y, D\.\[0-9]+\\);" 1 "gimple" { xfail *-*-* } } } */
/* { dg-final { scan-tree-dump-times "f \\(D\.\[0-9]+\\);" 1 "gimple" { xfail *-*-* } } } */

47
gcc/testsuite/c-c++-common/gomp/attrs-metadirective-1.c Normal file
View file

@ -0,0 +1,47 @@
/* { dg-do compile { target { c || c++11 } } } */
/* { dg-options "-fopenmp -std=c23" { target { c } } } */
#define N 100
void
f (int a[], int b[], int c[])
{
int i;
[[omp::directive (metadirective
default (teams loop)
default (parallel loop))]] /* { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective'" } */
for (i = 0; i < N; i++) c[i] = a[i] * b[i];
[[omp::directive (metadirective
default (bad_directive))]] /* { dg-error "unknown directive name before '\\)' token" } */
for (i = 0; i < N; i++) c[i] = a[i] * b[i];
[[omp::directive (metadirective
where (device={arch("nvptx")}: parallel loop) /* { dg-error "'where' is not valid for 'metadirective'" } */
default (teams loop))]]
for (i = 0; i < N; i++) c[i] = a[i] * b[i];
[[omp::directive (metadirective
otherwise (teams loop)
when (device={arch("nvptx")}: parallel loop))]] /* { dg-error "'otherwise' or 'default' clause must appear last" } */
for (i = 0; i < N; i++) c[i] = a[i] * b[i];
[[omp::directive (metadirective
when (device={arch("nvptx")} parallel loop) /* { dg-error "expected ':' before 'parallel'" } */
default (teams loop))]]
for (i = 0; i < N; i++) c[i] = a[i] * b[i];
[[omp::directive (metadirective
default (metadirective default (flush)))]] /* { dg-error "metadirectives cannot be used as variants of a 'metadirective' before 'default'" } */
for (i = 0; i < N; i++) c[i] = a[i] * b[i];
/* Test improperly nested metadirectives - even though the second
metadirective resolves to 'omp nothing', that is not the same as there
being literally nothing there. */
[[omp::directive (metadirective
when (implementation={vendor("gnu")}: parallel for))]]
[[omp::directive (metadirective /* { dg-error "loop nest expected" } */
when (implementation={vendor("cray")}: parallel for))]]
for (i = 0; i < N; i++) c[i] = a[i] * b[i];
}

76
gcc/testsuite/c-c++-common/gomp/attrs-metadirective-2.c Normal file
View file

@ -0,0 +1,76 @@
/* { dg-do compile { target { c || c++11 } } } */
/* { dg-options "-fopenmp -std=c23" { target { c } } } */
#define N 100
int
main (void)
{
int x = 0;
int y = 0;
/* Test implicit default (nothing). */
[[omp::directive (metadirective,
when (device={arch("nvptx")}: barrier))]]
x = 1;
/* Test with multiple standalone directives. */
[[omp::directive (metadirective,
when (device={arch("nvptx")}: barrier),
default (flush))]]
x = 1;
/* Test combining a standalone directive with one that takes a statement
body. */
[[omp::directive (metadirective,
when (device={arch("nvptx")}: parallel),
default (barrier))]]
x = 1;
/* Test combining a standalone directive with one that takes a for loop. */
[[omp::directive (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. */
[[omp::directive (metadirective,
when (device={arch("nvptx")}: parallel for),
default (parallel))]]
for (int i = 0; i < N; i++)
x += i;
/* Test labels inside statement body. */
[[omp::directive (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. */
[[omp::directive (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;
}

24
gcc/testsuite/c-c++-common/gomp/attrs-metadirective-3.c Normal file
View file

@ -0,0 +1,24 @@
/* { dg-do compile { target { c || c++11 } } } */
/* { dg-options "-fopenmp -std=c23" { target { c } } } */
/* { dg-additional-options "-fdump-tree-gimple" } */
#define N 100
void
f (int x[], int y[], int z[])
{
int i;
[[omp::sequence (directive (target map(to: x, y) map(from: z)),
directive (metadirective
when (device={arch("nvptx")}: teams loop)
default (parallel loop)))]]
for (i = 0; i < N; i++)
z[i] = x[i] * y[i];
}
/* If offload device "nvptx" isn't supported, the front end can eliminate
that alternative and not produce a metadirective at all. Otherwise this
won't be resolved until late. */
/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" { target { ! offload_nvptx } } } } */
/* { dg-final { scan-tree-dump "#pragma omp metadirective" "gimple" { target { offload_nvptx } } } } */

41
gcc/testsuite/c-c++-common/gomp/attrs-metadirective-4.c Normal file
View file

@ -0,0 +1,41 @@
/* { dg-do compile { target { c || c++11 } } } */
/* { dg-options "-fopenmp -std=c23" { target { c } } } */
/* { 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;
[[omp::directive (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 (void)
{
double a[N];
#pragma omp target teams map(from: a[0:N])
f (a, 3.14159);
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 "otherwise:" 1 "original" } } */
/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } } */
/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */

26
gcc/testsuite/c-c++-common/gomp/attrs-metadirective-5.c Normal file
View file

@ -0,0 +1,26 @@
/* { dg-do compile { target { c || c++11 } } } */
/* { dg-options "-fopenmp -std=c23" { target { c } } } */
/* { dg-additional-options "-fdump-tree-original" } */
#define N 100
void
f (int a[], int flag)
{
int i;
[[omp::directive (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" } } */

33
gcc/testsuite/c-c++-common/gomp/attrs-metadirective-6.c Normal file
View file

@ -0,0 +1,33 @@
/* { dg-do compile { target { c || c++11 } } } */
/* { dg-options "-fopenmp -std=c23" { target { c } } } */
/* { 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)
{
[[omp::directive (metadirective
when (user={condition(run_parallel)}: parallel))]]
{
int i;
[[omp::directive (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 "otherwise:" 2 "original" } } */
/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */

42
gcc/testsuite/c-c++-common/gomp/attrs-metadirective-7.c Normal file
View file

@ -0,0 +1,42 @@
/* { dg-do compile { target { c || c++11 } } } */
/* { dg-options "-fopenmp -std=c23" { target { c } } } */
/* { dg-additional-options "-fdump-tree-gimple -fdump-tree-ompdevlow" } */
#define N 256
void
f (int a[], int num)
{
int i;
[[omp::directive (metadirective
when (target_device={device_num(num), kind("gpu"), arch("nvptx")}:
target parallel for map(tofrom: a[0:N]))
when (target_device={device_num(num), kind("gpu"),
arch("amdgcn"), isa("gfx906")}:
target parallel for)
when (target_device={device_num(num), kind("cpu"), arch("x86_64")}:
parallel for))]]
for (i = 0; i < N; i++)
a[i] += i;
[[omp::directive (metadirective
when (target_device={kind("gpu"), arch("nvptx")}:
target parallel for map(tofrom: a[0:N])))]]
for (i = 0; i < N; i++)
a[i] += i;
}
/* For configurations with offloading, we expect one "pragma omp target"
with "device(num)" for each target_device selector that specifies
"device_num(num)". Without offloading, there should be zero as the
resolution happens during gimplification. */
/* { dg-final { scan-tree-dump-times "pragma omp target\[^\\n\]* device\\(num" 3 "gimple" { target offloading_enabled } } } */
/* { dg-final { scan-tree-dump-times "pragma omp target\[^\\n\]* device\\(num" 0 "gimple" { target { ! offloading_enabled } } } } */
/* For configurations with offloading, expect one OMP_TARGET_DEVICE_MATCHES
for each kind/arch/isa selector. These are supposed to go away after
ompdevlow. */
/* { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 9 "gimple" { target offloading_enabled } } } */
/* { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 0 "gimple" { target { ! offloading_enabled } } } } */
/* { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 0 "ompdevlow" { target offloading_enabled } } } */

18
gcc/testsuite/c-c++-common/gomp/attrs-metadirective-8.c Normal file
View file

@ -0,0 +1,18 @@
/* { dg-do compile { target { c || c++11 } } } */
/* { dg-options "-fopenmp -std=c23" { target { c } } } */
#define N 256
void
f (void)
{
int i;
int a[N];
[[omp::directive (metadirective
when( device={kind(nohost)}: nothing )
when( device={arch("nvptx")}: nothing)
default( parallel for))]]
for (i = 0; i < N; i++)
a[i] = i;
}

29
gcc/testsuite/c-c++-common/gomp/declare-variant-arg-exprs.c Normal file
View file

@ -0,0 +1,29 @@
/* { dg-do compile } */
/* { dg-additional-options "-foffload=disable" } */
/* { dg-additional-options "-mavx512bw -mavx512vl" { target { i?86-*-* x86_64-*-* } } } */
/* References to function parameters in dynamic selector expressions for
"declare variant" isn't supported yet; see PR 113904. Check to see that
a proper error is diagnosed meanwhile and GCC doesn't just wander off
into the weeds and ICE. */
extern int frob (int);
void f01 (int, int);
void f02 (int, int);
void f03 (int, int);
#pragma omp declare variant (f01) match (target_device={device_num (devnum), isa("avx512f","avx512vl")}) /* { dg-message "sorry, unimplemented: reference to function parameter" } */
#pragma omp declare variant (f02) match (implementation={vendor(score(15):gnu)})
#pragma omp declare variant (f03) match (user={condition(score(11):frob (ok + 42))}) /* { dg-message "sorry, unimplemented: reference to function parameter" } */
void f04 (int devnum, int ok);
void
test1 (void)
{
int i;
#pragma omp parallel for
for (i = 0; i < 1; i++)
f04 (17, 1);
}

26
gcc/testsuite/c-c++-common/gomp/declare-variant-dynamic-1.c Normal file
View file

@ -0,0 +1,26 @@
/* { dg-do compile } */
/* { dg-additional-options "-fdump-tree-gimple" } */
extern int foo_p (int);
extern int bar;
int f01 (int);
int f02 (int);
int f03 (int);
int f04 (int);
#pragma omp declare variant (f01) match (device={isa("avx512f")}) /* 4 */
#pragma omp declare variant (f02) match (implementation={vendor(score(3):gnu)},device={kind(cpu)}) /* 1 + 3 */
#pragma omp declare variant (f03) match (user={condition(score(9):foo_p (bar))})
#pragma omp declare variant (f04) match (implementation={vendor(score(6):gnu)},device={kind(host)}) /* 1 + 6 */
int f05 (int);
int
test1 (int x)
{
return f05 (x);
}
/* { dg-final { scan-tree-dump "f03 \\\(x" "gimple" } } */
/* { dg-final { scan-tree-dump "f04 \\\(x" "gimple" } } */
/* { dg-final { scan-tree-dump-not "f05 \\\(x" "gimple" } } */

30
gcc/testsuite/c-c++-common/gomp/declare-variant-dynamic-2.c Normal file
View file

@ -0,0 +1,30 @@
/* { dg-do compile } */
/* { dg-additional-options "-fdump-tree-gimple" } */
extern int foo_p (int);
extern int bar;
extern int omp_get_default_device (void);
int f01 (int);
int f02 (int);
int f03 (int);
int f04 (int);
#pragma omp declare variant (f01) match (target_device={device_num(omp_get_default_device()), isa("avx512f")}) /* 4 */
#pragma omp declare variant (f02) match (user={condition(score(6):0)})
#pragma omp declare variant (f03) match (user={condition(score(5):foo_p (bar))})
#pragma omp declare variant (f04) match (user={condition(score(3):0)})
int f05 (int);
int
test1 (int x)
{
return f05 (x);
}
/* f01 and f03 are the dynamic selectors, the fall-through is f05.
f02 and f04 are static selectors and do not match. */
/* { dg-final { scan-tree-dump "f01 \\\(x" "gimple" } } */
/* { dg-final { scan-tree-dump "f03 \\\(x" "gimple" } } */
/* { dg-final { scan-tree-dump "f05 \\\(x" "gimple" } } */
/* { dg-final { scan-tree-dump-not "f02 \\\(x" "gimple" } } */
/* { dg-final { scan-tree-dump-not "f04 \\\(x" "gimple" } } */

58
gcc/testsuite/c-c++-common/gomp/metadirective-1.c Normal file
View file

@ -0,0 +1,58 @@
/* { dg-do compile } */
#define N 100
void
f (int a[], int b[], int c[])
{
int i;
#pragma omp metadirective \
default (teams loop) \
default (parallel loop) /* { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective'" } */
for (i = 0; i < N; i++) c[i] = a[i] * b[i];
#pragma omp metadirective \
otherwise (teams loop) \
default (parallel loop) /* { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective'" } */
for (i = 0; i < N; i++) c[i] = a[i] * b[i];
#pragma omp metadirective \
otherwise (teams loop) \
otherwise (parallel loop) /* { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective'" } */
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 \
where (device={arch("nvptx")}: parallel loop) /* { dg-error "'where' is not valid for 'metadirective'" } */ \
default (teams loop)
for (i = 0; i < N; i++) c[i] = a[i] * b[i];
#pragma omp metadirective \
otherwise (teams loop) \
when (device={arch("nvptx")}: parallel loop) /* { dg-error "'otherwise' or 'default' clause must appear last" } */
for (i = 0; i < N; i++) c[i] = a[i] * b[i];
#pragma omp metadirective \
when (device={arch("nvptx")} parallel loop) /* { dg-error "expected ':' before 'parallel'" } */ \
default (teams loop)
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 variants of a 'metadirective' before 'default'" } */
for (i = 0; i < N; i++) c[i] = a[i] * b[i];
/* Test improperly nested metadirectives - even though the second
metadirective resolves to 'omp nothing', that is not the same as there
being literally nothing there. */
#pragma omp metadirective \
when (implementation={vendor("gnu")}: parallel for)
#pragma omp metadirective \
when (implementation={vendor("cray")}: parallel for)
/* { dg-error "loop nest expected before '#pragma'" "" { target c } .-2 } */
/* { dg-error "loop nest expected" "" { target c++ } .-3 } */
for (i = 0; i < N; i++) c[i] = a[i] * b[i];
}

75
gcc/testsuite/c-c++-common/gomp/metadirective-2.c Normal file
View file

@ -0,0 +1,75 @@
/* { 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;
}

23
gcc/testsuite/c-c++-common/gomp/metadirective-3.c Normal file
View file

@ -0,0 +1,23 @@
/* { dg-do compile } */
/* { dg-additional-options "-fdump-tree-gimple" } */
#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];
}
/* If offload device "nvptx" isn't supported, the front end can eliminate
that alternative and not produce a metadirective at all. Otherwise this
won't be resolved until late. */
/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" { target { ! offload_nvptx } } } } */
/* { dg-final { scan-tree-dump "#pragma omp metadirective" "gimple" { target { offload_nvptx } } } } */

40
gcc/testsuite/c-c++-common/gomp/metadirective-4.c Normal file
View file

@ -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 (void)
{
double a[N];
#pragma omp target teams map(from: a[0:N])
f (a, 3.14159);
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 "otherwise:" 1 "original" } } */
/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } } */
/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */

25
gcc/testsuite/c-c++-common/gomp/metadirective-5.c Normal file
View file

@ -0,0 +1,25 @@
/* { 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" } } */

32
gcc/testsuite/c-c++-common/gomp/metadirective-6.c Normal file
View file

@ -0,0 +1,32 @@
/* { 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 "otherwise:" 2 "original" } } */
/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */

41
gcc/testsuite/c-c++-common/gomp/metadirective-7.c Normal file
View file

@ -0,0 +1,41 @@
/* { dg-do compile } */
/* { dg-additional-options "-fdump-tree-gimple -fdump-tree-ompdevlow" } */
#define N 256
void
f (int a[], int num)
{
int i;
#pragma omp metadirective \
when (target_device={device_num(num), kind("gpu"), arch("nvptx")}: \
target parallel for map(tofrom: a[0:N])) \
when (target_device={device_num(num), kind("gpu"), \
arch("amdgcn"), isa("gfx906")}: \
target parallel for) \
when (target_device={device_num(num), kind("cpu"), arch("x86_64")}: \
parallel for)
for (i = 0; i < N; i++)
a[i] += i;
#pragma omp metadirective \
when (target_device={kind("gpu"), arch("nvptx")}: \
target parallel for map(tofrom: a[0:N]))
for (i = 0; i < N; i++)
a[i] += i;
}
/* For configurations with offloading, we expect one "pragma omp target"
with "device(num)" for each target_device selector that specifies
"device_num(num)". Without offloading, there should be zero as the
resolution happens during gimplification. */
/* { dg-final { scan-tree-dump-times "pragma omp target\[^\\n\]* device\\(num" 3 "gimple" { target offloading_enabled } } } */
/* { dg-final { scan-tree-dump-times "pragma omp target\[^\\n\]* device\\(num" 0 "gimple" { target { ! offloading_enabled } } } } */
/* For configurations with offloading, expect one OMP_TARGET_DEVICE_MATCHES
for each kind/arch/isa selector. These are supposed to go away after
ompdevlow. */
/* { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 9 "gimple" { target offloading_enabled } } } */
/* { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 0 "gimple" { target { ! offloading_enabled } } } } */
/* { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 0 "ompdevlow" { target offloading_enabled } } } */

17
gcc/testsuite/c-c++-common/gomp/metadirective-8.c Normal file
View file

@ -0,0 +1,17 @@
/* { dg-do compile } */
#define N 256
void
f (void)
{
int i;
int a[N];
#pragma omp metadirective \
when( device={kind(nohost)}: nothing ) \
when( device={arch("nvptx")}: nothing) \
default( parallel for)
for (i = 0; i < N; i++)
a[i] = i;
}

178
gcc/testsuite/c-c++-common/gomp/metadirective-construct.c Normal file
View file

@ -0,0 +1,178 @@
/* { dg-do compile } */
/* { dg-additional-options "-foffload=disable -fdump-tree-optimized" } */
#include <stdlib.h>
static void
init (int n, double *a)
{
for (int i = 0; i < n; i++)
a[i] = (double) i;
}
static void
check (int n, double *a, double s)
{
for (int i = 0; i < n; i++)
if (a[i] != (double) i * s)
abort ();
}
typedef void (transform_fn) (int, double *, double);
static void doit (transform_fn *f, int n, double *a, double s)
{
init (n, a);
(*f) (n, a, s);
check (n, a, s);
}
/* Check various combinations for enforcing correct ordering of
construct matches. */
static void
f1 (int n, double* a, double s)
{
#pragma omp target teams
#pragma omp parallel
#pragma omp metadirective \
when (construct={target} \
: for) \
default (error at(execution) message("f1 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
static void
f2 (int n, double* a, double s)
{
#pragma omp target teams
#pragma omp parallel
#pragma omp metadirective \
when (construct={teams, parallel} \
: for) \
default (error at(execution) message("f2 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
static void
f3 (int n, double* a, double s)
{
#pragma omp target teams
#pragma omp parallel
#pragma omp metadirective \
when (construct={target, teams, parallel} \
: for) \
default (error at(execution) message("f3 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
static void
f4 (int n, double* a, double s)
{
#pragma omp target teams
#pragma omp parallel
#pragma omp metadirective \
when (construct={target, parallel} \
: for) \
default (error at(execution) message("f4 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
static void
f5 (int n, double* a, double s)
{
#pragma omp target teams
#pragma omp parallel
#pragma omp metadirective \
when (construct={target, teams} \
: for) \
default (error at(execution) message("f5 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
/* Next batch is for things where the construct doesn't match the context. */
static void
f6 (int n, double* a, double s)
{
#pragma omp target
#pragma omp teams
#pragma omp metadirective \
when (construct={parallel} \
: error at(execution) message("f6 match failed")) \
default (parallel for)
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
static void
f7 (int n, double* a, double s)
{
#pragma omp target
#pragma omp teams
#pragma omp metadirective \
when (construct={target, parallel} \
: error at(execution) message("f7 match failed")) \
default (parallel for)
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
static void
f8 (int n, double* a, double s)
{
#pragma omp target
#pragma omp teams
#pragma omp metadirective \
when (construct={parallel, target} \
: error at(execution) message("f8 match failed")) \
default (parallel for)
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
/* Next test choosing the best alternative when there are multiple
matches. */
static void
f9 (int n, double* a, double s)
{
#pragma omp target teams
#pragma omp parallel
#pragma omp metadirective \
when (construct={teams, parallel} \
: error at(execution) message("f9 match incorrect 1")) \
when (construct={target, teams, parallel} \
: for) \
when (construct={target, teams} \
: error at(execution) message("f9 match incorrect 2")) \
default (error at(execution) message("f9 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
/* Note there are no tests for the matching the extended simd clause
syntax, which is only useful for "declare variant". */
#define N 10
#define S 2.0
int
main (void)
{
double a[N];
doit (f1, N, a, S);
doit (f2, N, a, S);
doit (f3, N, a, S);
doit (f4, N, a, S);
doit (f5, N, a, S);
doit (f6, N, a, S);
doit (f7, N, a, S);
doit (f8, N, a, S);
doit (f9, N, a, S);
}
/* All the error calls should be optimized away. */
/* { dg-final { scan-tree-dump-not "__builtin_GOMP_error" "optimized" } } */

149
gcc/testsuite/c-c++-common/gomp/metadirective-device.c Normal file
View file

@ -0,0 +1,149 @@
/* { dg-do compile } */
/* { dg-additional-options "-foffload=disable -fdump-tree-optimized" } */
/* { dg-additional-options "-DDEVICE_ARCH=x86_64 -DDEVICE_ISA=sse -msse" { target x86_64-*-* } } */
#include <stdlib.h>
static void
init (int n, double *a)
{
for (int i = 0; i < n; i++)
a[i] = (double) i;
}
static void
check (int n, double *a, double s)
{
for (int i = 0; i < n; i++)
if (a[i] != (double) i * s)
abort ();
}
typedef void (transform_fn) (int, double *, double);
static void
doit (transform_fn *f, int n, double *a, double s)
{
init (n, a);
(*f) (n, a, s);
check (n, a, s);
}
/* Check kind=host matches (with offloading disabled). */
static void
f1 (int n, double* a, double s)
{
#pragma omp metadirective \
when (device={kind(host)} \
: parallel for) \
default (error at(execution) message("f1 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
/* Check kind=nohost does not match (with offloading disabled). */
static void
f2 (int n, double* a, double s)
{
#pragma omp metadirective \
when (device={kind(nohost)} \
: error at(execution) message("f2 match failed")) \
default (parallel for)
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
/* Check arch. Either DEVICE_ARCH is defined by command-line option,
or we know it is not x86_64. */
#ifdef DEVICE_ARCH
static void
f3 (int n, double* a, double s)
{
#pragma omp metadirective \
when (device={arch(DEVICE_ARCH)} \
: parallel for) \
default (error at(execution) message("f3 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
#else
static void
f3 (int n, double* a, double s)
{
#pragma omp metadirective \
when (device={arch("x86_64")} \
: error at(execution) message("f3 match failed")) \
default (parallel for)
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
#endif
/* Check both kind and arch together. */
#ifdef DEVICE_ARCH
static void
f4 (int n, double* a, double s)
{
#pragma omp metadirective \
when (device={arch(DEVICE_ARCH), kind(host)} \
: parallel for) \
default (error at(execution) message("f4 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
#else
static void
f4 (int n, double* a, double s)
{
#pragma omp metadirective \
when (device={arch("x86_64"), kind(host)} \
: error at(execution) message("f4 match failed")) \
default (parallel for)
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
#endif
/* Check kind, arch, and ISA together. */
#if defined(DEVICE_ARCH) && defined(DEVICE_ISA)
static void
f5 (int n, double* a, double s)
{
#pragma omp metadirective \
when (device={arch(DEVICE_ARCH), kind(host), isa(DEVICE_ISA)} \
: parallel for) \
default (error at(execution) message("f5 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
#else
static void
f5 (int n, double* a, double s)
{
#pragma omp metadirective \
when (device={arch("x86_64"), kind(host), isa("sse")} \
: error at(execution) message("f5 match failed")) \
default (parallel for)
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
#endif
#define N 10
#define S 2.0
int
main (void)
{
double a[N];
doit (f1, N, a, S);
doit (f2, N, a, S);
doit (f3, N, a, S);
doit (f4, N, a, S);
doit (f5, N, a, S);
}
/* All the metadirectives involving the device selector should be
fully resolved and the error calls optimized away. */
/* { dg-final { scan-tree-dump-not "__builtin_GOMP_error" "optimized" } } */

95
gcc/testsuite/c-c++-common/gomp/metadirective-no-score.c Normal file
View file

@ -0,0 +1,95 @@
/* { dg-do compile { target x86_64-*-* } } */
/* { dg-additional-options "-foffload=disable" } */
/* This test is expected to fail with compile-time errors:
"A trait-score cannot be specified in traits from the construct,
device or target_device trait-selector-sets." */
/* Define this to avoid dependence on libgomp header files. */
#define omp_initial_device -1
void
f1 (int n, double *a, double s)
{
#pragma omp metadirective \
when (device={kind (score(5) : host)} \
: parallel for)
/* { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-2 } */
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
void
f2 (int n, double *a, double s)
{
#pragma omp metadirective \
when (device={kind (host), arch (score(6) : x86_64), isa (avx512f)} \
: parallel for)
/* { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-2 } */
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
void
f3 (int n, double *a, double s)
{
#pragma omp metadirective \
when (device={kind (host), arch (score(6) : x86_64), \
isa (score(7): avx512f)} \
: parallel for)
/* { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-3 } */
/* { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-3 } */
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
void
f4 (int n, double *a, double s)
{
#pragma omp metadirective \
when (target_device={device_num (score(42) : omp_initial_device), \
kind (host)} \
: parallel for)
/* { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-3 } */
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
void
f5 (int n, double *a, double s)
{
#pragma omp metadirective \
when (target_device={device_num(omp_initial_device), \
kind (score(5) : host)} \
: parallel for)
/* { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-2 } */
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
void
f6 (int n, double *a, double s)
{
#pragma omp metadirective \
when (target_device={device_num(omp_initial_device), kind (host), \
arch (score(6) : x86_64), isa (avx512f)} \
: parallel for)
/* { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-2 } */
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
void
f7 (int n, double *a, double s)
{
#pragma omp metadirective \
when (target_device={device_num(omp_initial_device), kind (host), \
arch (score(6) : x86_64), \
isa (score(7): avx512f)} \
: parallel for)
/* { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-3 } */
/* { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-3 } */
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}

151
gcc/testsuite/c-c++-common/gomp/metadirective-target-device-1.c Normal file
View file

@ -0,0 +1,151 @@
/* { dg-do compile } */
/* { dg-additional-options "-fdump-tree-optimized" } */
/* { dg-additional-options "-DDEVICE_ARCH=x86_64 -DDEVICE_ISA=mmx -mmmx" { target x86_64-*-* } } */
#include <stdlib.h>
static void
init (int n, double *a)
{
for (int i = 0; i < n; i++)
a[i] = (double) i;
}
static void
check (int n, double *a, double s)
{
for (int i = 0; i < n; i++)
if (a[i] != (double) i * s)
abort ();
}
typedef void (transform_fn) (int, double *, double);
static void
doit (transform_fn *f, int n, double *a, double s)
{
init (n, a);
(*f) (n, a, s);
check (n, a, s);
}
/* Check kind=host matches (with offloading disabled). */
static void
f1 (int n, double* a, double s)
{
#pragma omp metadirective \
when (target_device={kind(host)} \
: parallel for) \
default (error at(execution) message("f1 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
/* Check kind=nohost does not match (with offloading disabled). */
static void
f2 (int n, double* a, double s)
{
#pragma omp metadirective \
when (target_device={kind(nohost)} \
: error at(execution) message("f2 match failed")) \
default (parallel for)
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
/* Check arch. Either DEVICE_ARCH is defined by command-line option,
or we know it is not x86_64. */
#ifdef DEVICE_ARCH
static void
f3 (int n, double* a, double s)
{
#pragma omp metadirective \
when (target_device={arch(DEVICE_ARCH)} \
: parallel for) \
default (error at(execution) message("f3 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
#else
static void
f3 (int n, double* a, double s)
{
#pragma omp metadirective \
when (target_device={arch("x86_64")} \
: error at(execution) message("f3 match failed")) \
default (parallel for)
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
#endif
/* Check both kind and arch together. */
#ifdef DEVICE_ARCH
static void
f4 (int n, double* a, double s)
{
#pragma omp metadirective \
when (target_device={arch(DEVICE_ARCH), kind(host)} \
: parallel for) \
default (error at(execution) message("f4 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
#else
static void
f4 (int n, double* a, double s)
{
#pragma omp metadirective \
when (target_device={arch("x86_64"), kind(host)} \
: error at(execution) message("f4 match failed")) \
default (parallel for)
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
#endif
/* Check kind, arch, and ISA together. */
#if defined(DEVICE_ARCH) && defined(DEVICE_ISA)
static void
f5 (int n, double* a, double s)
{
#pragma omp metadirective \
when (target_device={arch(DEVICE_ARCH), kind(host), isa(DEVICE_ISA)} \
: parallel for) \
default (error at(execution) message("f5 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
#else
static void
f5 (int n, double* a, double s)
{
#pragma omp metadirective \
when (target_device={arch("x86_64"), kind(host), isa("mmx")} \
: error at(execution) message("f5 match failed")) \
default (parallel for)
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
#endif
#define N 10
#define S 2.0
int
main (void)
{
double a[N];
doit (f1, N, a, S);
doit (f2, N, a, S);
doit (f3, N, a, S);
doit (f4, N, a, S);
doit (f5, N, a, S);
}
/* The gimplifier emits all 5 alternatives. In configurations without
offloading, the dynamic selector can be completely resolved, otherwise
the calls to GOMP_error must remain in the generated code pending
runtime evaluation of the selectors. */
/* { dg-final { scan-tree-dump-times "GOMP_error" 5 "optimized" { target offloading_enabled } } } */
/* { dg-final { scan-tree-dump-times "GOMP_error" 0 "optimized" { target { ! offloading_enabled } } } } */

132
gcc/testsuite/c-c++-common/gomp/metadirective-target-device-2.c Normal file
View file

@ -0,0 +1,132 @@
/* { dg-do compile */
/* { dg-additional-options "-fdump-tree-optimized" } */
/* In configurations without offloading configured, we can resolve many
instances of the target_device context selector at gimplification time
instead of waiting until late resolution. */
/* Device 0 may be either the host or an offloading device, in configurations
that support them. */
void
f1 (int n, double* a, double s)
{
#pragma omp metadirective \
when (target_device={device_num(0), kind(host)} \
: parallel for) \
default (error at(execution) message("f1 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
/* { dg-final { scan-tree-dump-not "f1 match failed" "optimized" { target { ! offloading_enabled } } } } */
/* { dg-final { scan-tree-dump "f1 match failed" "optimized" { target offloading_enabled } } } */
/* Device -1 is always the host, even in offloading configurations. */
void
f2 (int n, double* a, double s)
{
#pragma omp metadirective \
when (target_device={device_num(-1), kind(host)} \
: parallel for) \
default (error at(execution) message("f2 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
/* { dg-final { scan-tree-dump-not "f2 match failed" "optimized" } } */
/* Constant device numbers < -1 cause the match to fail, even in offloading
configurations. */
void
f3 (int n, double* a, double s)
{
#pragma omp metadirective \
when (target_device={device_num(-42), kind(host)} \
: error at(execution) message("f3 match failed")) \
default (parallel for)
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
/* { dg-final { scan-tree-dump-not "f3 match failed" "optimized" } } */
/* In a non-offloading configuration, devices > 0 are always invalid and
will cause the match to fail. In an offload configuration, we don't
know at compile-time if the device number is valid or if it refers to
a host or offload device. */
void
f4 (int n, double* a, double s)
{
#pragma omp metadirective \
when (target_device={device_num(42), kind(host)} \
: error at(execution) message("f4 match failed")) \
default (parallel for)
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
/* { dg-final { scan-tree-dump-not "f4 match failed" "optimized" { target { ! offloading_enabled } } } } */
/* { dg-final { scan-tree-dump "f4 match failed" "optimized" { target offloading_enabled } } } */
extern int omp_get_initial_device (void);
extern int omp_get_default_device (void);
extern int omp_get_num_devices (void);
extern int omp_get_device_num (void);
/* On a non-offloading configuration, omp_get_default_device() always refers
to the host. With offloading, we know it's a valid device number but
not whether it's the host. */
void
f5 (int n, double* a, double s)
{
#pragma omp metadirective \
when (target_device={device_num(omp_get_default_device ()), kind(host)} \
: parallel for) \
default (error at(execution) message("f5 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
/* { dg-final { scan-tree-dump-not "f5 match failed" "optimized" { target { ! offloading_enabled } } } } */
/* { dg-final { scan-tree-dump "f5 match failed" "optimized" { target offloading_enabled } } } */
/* omp_get_initial_device() always refers to the host, whether offloading is
configured or not. */
void
f6 (int n, double* a, double s)
{
#pragma omp metadirective \
when (target_device={device_num(omp_get_initial_device ()), kind(host)} \
: parallel for) \
default (error at(execution) message("f6 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
/* { dg-final { scan-tree-dump-not "f6 match failed" "optimized" } } */
/* omp_get_num_devices() returns the number of non-host offload devices,
therefore using it as a device number always refers to the host, whether
offloading is enabled or not. */
void
f7 (int n, double* a, double s)
{
#pragma omp metadirective \
when (target_device={device_num(omp_get_num_devices ()), kind(host)} \
: parallel for) \
default (error at(execution) message("f7 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
/* { dg-final { scan-tree-dump-not "f7 match failed" "optimized" } } */
/* omp_get_device_num() always returns a valid device number but we don't
know whether it's the host or offload device, in configurations that
support offloading. */
void
f8 (int n, double* a, double s)
{
#pragma omp metadirective \
when (target_device={device_num(omp_get_device_num ()), kind(host)} \
: parallel for) \
default (error at(execution) message("f8 match failed"))
for (int i = 0; i < n; i++)
a[i] = a[i] * s;
}
/* { dg-final { scan-tree-dump-not "f8 match failed" "optimized" { target { ! offloading_enabled } } } } */
/* { dg-final { scan-tree-dump "f8 match failed" "optimized" { target offloading_enabled } } } */

37
libgomp/testsuite/libgomp.c-c++-common/metadirective-1.c Normal file
View file

@ -0,0 +1,37 @@
/* { 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;
}

41
libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c Normal file
View file

@ -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 (void)
{
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;
f (a, M_E);
for (i = 0; i < N; i++)
if (fabs (a[i] - (M_E * i)) > EPSILON)
return 1;
return 0;
}

36
libgomp/testsuite/libgomp.c-c++-common/metadirective-3.c Normal file
View file

@ -0,0 +1,36 @@
/* { 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;
}

54
libgomp/testsuite/libgomp.c-c++-common/metadirective-4.c Normal file
View file

@ -0,0 +1,54 @@
/* { 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;
}

48
libgomp/testsuite/libgomp.c-c++-common/metadirective-5.c Normal file
View file

@ -0,0 +1,48 @@
/* { dg-do run } */
#define N 100
#include <stdio.h>
#include <omp.h>
int
f (int a[], int num)
{
int on_device = 0;
int i;
#pragma omp metadirective \
when (target_device={device_num(num), kind("gpu")}: \
target parallel for map(to: a[0:N]), map(from: on_device)) \
default (parallel for private (on_device))
for (i = 0; i < N; i++)
{
a[i] += i;
on_device = 1;
}
return on_device;
}
int
main (void)
{
int a[N];
int on_device_count = 0;
int i;
for (i = 0; i < N; i++)
a[i] = i;
for (i = 0; i <= omp_get_num_devices (); i++)
on_device_count += f (a, i);
if (on_device_count != omp_get_num_devices ())
return 1;
for (i = 0; i < N; i++)
if (a[i] != 2 * i)
return 2;
return 0;
}

66
libgomp/testsuite/libgomp.c-c++-common/metadirective-late-1.c Normal file
View file

@ -0,0 +1,66 @@
/* { dg-do run } */
/* { dg-additional-options -O0 } */
/* Test late resolution of metadirectives with dynamic selectors
in "declare simd" functions. All the variants do the same thing;
the purpose of this test is to ensure that the "condition" predicates
are all called, and in the correct order. */
static int pcount = 0;
static int __attribute__ ((noinline))
ptrue (int n)
{
pcount++;
if (pcount != n)
__builtin_abort ();
return 1;
}
static int __attribute__ ((noinline))
pfalse (int n)
{
pcount++;
if (pcount != n)
__builtin_abort ();
return 0;
}
#define N 256
#pragma omp declare simd
void
f (int a[])
{
int i;
#pragma omp metadirective \
when (construct={simd}: \
nothing) \
when (user={condition(score (100): pfalse (1))}: \
nothing) \
when (user={condition(score (90): pfalse (2))}: \
nothing) \
when (user={condition(score (70): (ptrue (5) && pfalse (6)))}: \
nothing) \
when (user={condition(score (80): (pfalse (3) || pfalse (4)))}: \
nothing) \
when (user={condition(score (60): \
(ptrue (7) ? pfalse (8) : ptrue (8)))}: \
nothing) \
otherwise (nothing)
for (i = 0; i < N; i++)
a[i] += i;
}
int a[N];
int
main (void)
{
f (a);
for (int i = 0; i < N; i++)
if (a[i] != i)
return 1;
return 0;
}

66
libgomp/testsuite/libgomp.c-c++-common/metadirective-late-2.c Normal file
View file

@ -0,0 +1,66 @@
/* { dg-do run } */
/* { dg-additional-options -O3 } */
/* Test late resolution of metadirectives with dynamic selectors
in "declare simd" functions. All the variants do the same thing;
the purpose of this test is to ensure that the "condition" predicates
are all called, and in the correct order. */
static int pcount = 0;
static int __attribute__ ((noinline))
ptrue (int n)
{
pcount++;
if (pcount != n)
__builtin_abort ();
return 1;
}
static int __attribute__ ((noinline))
pfalse (int n)
{
pcount++;
if (pcount != n)
__builtin_abort ();
return 0;
}
#define N 256
#pragma omp declare simd
void
f (int a[])
{
int i;
#pragma omp metadirective \
when (construct={simd}: \
nothing) \
when (user={condition(score (100): pfalse (1))}: \
nothing) \
when (user={condition(score (90): pfalse (2))}: \
nothing) \
when (user={condition(score (70): (ptrue (5) && pfalse (6)))}: \
nothing) \
when (user={condition(score (80): (pfalse (3) || pfalse (4)))}: \
nothing) \
when (user={condition(score (60): \
(ptrue (7) ? pfalse (8) : ptrue (8)))}: \
nothing) \
otherwise (nothing)
for (i = 0; i < N; i++)
a[i] += i;
}
int a[N];
int
main (void)
{
f (a);
for (int i = 0; i < N; i++)
if (a[i] != i)
return 1;
return 0;
}

76
libgomp/testsuite/libgomp.c-c++-common/metadirective-target-device.c Normal file
View file

@ -0,0 +1,76 @@
#include <omp.h>
#include <stdlib.h>
#include <stdio.h>
/* PR112779 item (B) */
/* Check that the target_device selector correctly matches device numbers
and handles kind=host|nohost|any. */
static int
check_explicit_device (int d, int expect_host)
{
int ok = 0;
if (expect_host)
{
#pragma omp metadirective \
when (target_device={device_num(d), kind("host")} : nothing) \
otherwise (error at(execution) message("check_explicit_device host"))
ok = 1;
}
else
{
#pragma omp metadirective \
when (target_device={device_num(d), kind("nohost")} : nothing) \
otherwise (error at(execution) message("check_explicit_device nohost"))
ok = 1;
}
return ok;
}
static int
check_implicit_device (int d, int expect_host)
{
int ok = 0;
omp_set_default_device (d);
if (expect_host)
{
#pragma omp metadirective \
when (target_device={kind("host")} : nothing) \
otherwise (error at(execution) message("check_implicit_device host"))
ok = 1;
}
else
{
#pragma omp metadirective \
when (target_device={kind("nohost")} : nothing) \
otherwise (error at(execution) message("check_implicit_device nohost"))
ok = 1;
}
#pragma omp metadirective \
when (target_device={kind("any")} : nothing) \
otherwise (error at(execution) message("check_implicit_device any"))
ok = 1;
omp_set_default_device (omp_initial_device);
return ok;
}
int
main (void)
{
printf ("Checking omp_initial_device\n");
check_explicit_device (omp_initial_device, 1);
check_implicit_device (omp_initial_device, 1);
int n = omp_get_num_devices ();
printf ("There are %d devices\n", n);
for (int i = 0; i < n; i++)
{
printf ("Checking device %d\n", i);
check_explicit_device (i, 0);
check_implicit_device (i, 0);
}
return 0;
}