mirror of
git://gcc.gnu.org/git/gcc.git
synced 2025-02-17 22:30:31 +08:00
OpenMP: common C/C++ testcases for dispatch + adjust_args
gcc/testsuite/ChangeLog: * c-c++-common/gomp/declare-variant-2.c: Adjust dg-error directives. * c-c++-common/gomp/adjust-args-1.c: New test. * c-c++-common/gomp/adjust-args-2.c: New test. * c-c++-common/gomp/declare-variant-dup-match-clause.c: New test. * c-c++-common/gomp/dispatch-1.c: New test. * c-c++-common/gomp/dispatch-2.c: New test. * c-c++-common/gomp/dispatch-3.c: New test. * c-c++-common/gomp/dispatch-4.c: New test. * c-c++-common/gomp/dispatch-5.c: New test. * c-c++-common/gomp/dispatch-6.c: New test. * c-c++-common/gomp/dispatch-7.c: New test. * c-c++-common/gomp/dispatch-8.c: New test. * c-c++-common/gomp/dispatch-9.c: New test. * c-c++-common/gomp/dispatch-10.c: New test. libgomp/ChangeLog: * testsuite/libgomp.c-c++-common/dispatch-1.c: New test. * testsuite/libgomp.c-c++-common/dispatch-2.c: New test.
This commit is contained in:
Paul-Antoine Arras
parent
ed49709acd
commit
377eff7c38
30
gcc/testsuite/c-c++-common/gomp/adjust-args-1.c
Normal file
30
gcc/testsuite/c-c++-common/gomp/adjust-args-1.c
Normal file
@ -0,0 +1,30 @@
|
||||
/* { dg-do compile } */
|
||||
/* { dg-additional-options "-fdump-tree-gimple" } */
|
||||
|
||||
int f (int a, void *b, float c[2]);
|
||||
|
||||
#pragma omp declare variant (f) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: b, c)
|
||||
int f0 (int a, void *b, float c[2]);
|
||||
#pragma omp declare variant (f) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: b) adjust_args (need_device_ptr: c)
|
||||
int f1 (int a, void *b, float c[2]);
|
||||
|
||||
int test () {
|
||||
int a;
|
||||
void *b;
|
||||
float c[2];
|
||||
struct {int a;} s;
|
||||
|
||||
s.a = f0 (a, b, c);
|
||||
#pragma omp dispatch
|
||||
s.a = f0 (a, b, c);
|
||||
|
||||
f1 (a, b, c);
|
||||
#pragma omp dispatch
|
||||
s.a = f1 (a, b, c);
|
||||
|
||||
return s.a;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device \\(\\);" 2 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&c, D\.\[0-9]+\\);" 2 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(b, D\.\[0-9]+\\);" 2 "gimple" } } */
|
||||
31
gcc/testsuite/c-c++-common/gomp/adjust-args-2.c
Normal file
31
gcc/testsuite/c-c++-common/gomp/adjust-args-2.c
Normal file
@ -0,0 +1,31 @@
|
||||
/* { dg-do compile } */
|
||||
/* { dg-additional-options "-fdump-tree-gimple" } */
|
||||
|
||||
int f (int a, void *b, float c[2]);
|
||||
|
||||
#pragma omp declare variant (f) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: b, c)
|
||||
int f0 (int a, void *b, float c[2]);
|
||||
#pragma omp declare variant (f) adjust_args (need_device_ptr: b, c) match (construct={dispatch}) adjust_args (nothing: a)
|
||||
int f1 (int a, void *b, float c[2]);
|
||||
|
||||
void test () {
|
||||
int a;
|
||||
void *b;
|
||||
float c[2];
|
||||
|
||||
#pragma omp dispatch
|
||||
f0 (a, b, c);
|
||||
|
||||
#pragma omp dispatch device (-4852)
|
||||
f0 (a, b, c);
|
||||
|
||||
#pragma omp dispatch device (a + a)
|
||||
f0 (a, b, c);
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device \\(\\);" 3 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&c, D\.\[0-9]+\\);" 2 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(b, D\.\[0-9]+\\);" 2 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&c, -4852\\);" 1 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(b, -4852\\);" 1 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-not "#pragma omp dispatch device" "gimple" } } */
|
||||
@ -8,9 +8,9 @@ void f3 (void);
|
||||
void f4 (void);
|
||||
#pragma omp declare variant match(user={condition(0)}) /* { dg-error "expected '\\(' before 'match'" } */
|
||||
void f5 (void);
|
||||
#pragma omp declare variant (f1) /* { dg-error "expected 'match' before end of line" } */
|
||||
#pragma omp declare variant (f1) /* { dg-error "expected 'match' clause before end of line" } */
|
||||
void f6 (void);
|
||||
#pragma omp declare variant (f1) simd /* { dg-error "expected 'match' before 'simd'" } */
|
||||
#pragma omp declare variant (f1) simd /* { dg-error "expected 'match' clause before 'simd'" } */
|
||||
void f7 (void);
|
||||
#pragma omp declare variant (f1) match /* { dg-error "expected '\\(' before end of line" } */
|
||||
void f8 (void);
|
||||
|
||||
@ -0,0 +1,6 @@
|
||||
/* Check that errors are detected if a match clause appears more than once,
|
||||
even if compatible. */
|
||||
|
||||
void f();
|
||||
#pragma omp declare variant(f) match(construct={target}) match(construct={target}) /* { dg-error "too many 'match' clauses" } */
|
||||
void g();
|
||||
71
gcc/testsuite/c-c++-common/gomp/dispatch-1.c
Normal file
71
gcc/testsuite/c-c++-common/gomp/dispatch-1.c
Normal file
@ -0,0 +1,71 @@
|
||||
#include <stdint.h>
|
||||
|
||||
int f0 (int, long, double);
|
||||
void f2 (void);
|
||||
int f3 (void);
|
||||
void (*f4) (void);
|
||||
void f5 (int*, int);
|
||||
|
||||
int *pp (void);
|
||||
int dd (void);
|
||||
|
||||
void f1 (void)
|
||||
{
|
||||
int a, c;
|
||||
long b;
|
||||
double x;
|
||||
struct {int a; float b; short c;} s, *sp;
|
||||
int arr[3];
|
||||
|
||||
#pragma omp dispatch
|
||||
c = f0 (a, b, x);
|
||||
#pragma omp dispatch
|
||||
x = f0 (a * 4, 2 - b, x * x);
|
||||
#pragma omp dispatch
|
||||
s.a = f0 (a, sp->c, x);
|
||||
#pragma omp dispatch
|
||||
sp->c = f0 (s.a - 2, b / 3, x * 5);
|
||||
#pragma omp dispatch
|
||||
arr[0] = f0 (arr[1], !b, arr[2]);
|
||||
#pragma omp dispatch
|
||||
(*sp).c = f0 (s.a, b, x);
|
||||
#pragma omp dispatch
|
||||
sp->b = f0 (s.a++, b % 4, --x);
|
||||
#pragma omp dispatch
|
||||
f0 (f3(), b, s.b);
|
||||
#pragma omp dispatch
|
||||
f2 ();
|
||||
#pragma omp dispatch
|
||||
f4 ();
|
||||
#pragma omp dispatch
|
||||
f5 (pp(), pp()[dd()]);
|
||||
|
||||
#pragma omp dispatch nocontext(sp->a * x + arr[2])
|
||||
f2 ();
|
||||
#pragma omp dispatch nocontext(arr - (intptr_t)(x / s.b))
|
||||
f2 ();
|
||||
#pragma omp dispatch nocontext(x == s.c || b != c)
|
||||
f2 ();
|
||||
#pragma omp dispatch novariants(b << sp->c)
|
||||
f2 ();
|
||||
#pragma omp dispatch novariants(!arr | s.a)
|
||||
f2 ();
|
||||
#pragma omp dispatch novariants(s.c ? f3() : a & c)
|
||||
f2 ();
|
||||
#pragma omp dispatch nowait
|
||||
f2 ();
|
||||
#pragma omp dispatch device(-25373654)
|
||||
f2 ();
|
||||
#pragma omp dispatch device(b * (int)(x - sp->b))
|
||||
f2 ();
|
||||
#pragma omp dispatch is_device_ptr(arr)
|
||||
f2 ();
|
||||
#pragma omp dispatch is_device_ptr(sp)
|
||||
f2 ();
|
||||
#pragma omp dispatch depend(inout: sp)
|
||||
f2 ();
|
||||
#pragma omp dispatch depend(inoutset: arr[:2])
|
||||
f2 ();
|
||||
#pragma omp dispatch depend(out: arr)
|
||||
f2 ();
|
||||
}
|
||||
36
gcc/testsuite/c-c++-common/gomp/dispatch-10.c
Normal file
36
gcc/testsuite/c-c++-common/gomp/dispatch-10.c
Normal file
@ -0,0 +1,36 @@
|
||||
/* { dg-do compile } */
|
||||
/* { dg-additional-options "-fdump-tree-gimple" } */
|
||||
|
||||
int *f();
|
||||
|
||||
struct t {
|
||||
int *a, *b;
|
||||
};
|
||||
|
||||
void construct(int *x, int *y);
|
||||
void noconstruct(int *x, int *y);
|
||||
|
||||
#pragma omp declare variant(construct) match(construct={dispatch}) adjust_args(need_device_ptr: x,y)
|
||||
#pragma omp declare variant(noconstruct) match(implementation={vendor(gnu)})
|
||||
void bar(int *x, int *y);
|
||||
|
||||
int nocontext, novariant;
|
||||
|
||||
void sub(struct t *s, void *y)
|
||||
{
|
||||
bar( f(), s->b);
|
||||
#pragma omp dispatch device(0) is_device_ptr(s)
|
||||
bar( f(), s->b);
|
||||
|
||||
bar( (int *) y, s->b);
|
||||
#pragma omp dispatch device(0) is_device_ptr(y)
|
||||
bar( (int *) y, s->b);
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device" 2 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "__builtin_omp_set_default_device" 4 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "__builtin_omp_set_default_device \\(0\\);" 2 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "__builtin_omp_set_default_device \\(D\\.\[0-9\]+\\);" 2 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "__builtin_omp_get_mapped_ptr" 3 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "D\\.\[0-9\]+ = s->b;\[\r\n\]* *D\\.\[0-9\]+ = __builtin_omp_get_mapped_ptr \\(D\\.\[0-9\]+, 0\\);" 2 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "D\\.\[0-9\]+ = f \\(\\);\[\r\n\]* *D\\.\[0-9\]+ = __builtin_omp_get_mapped_ptr \\(D\\.\[0-9\]+, 0\\);" 1 "gimple" } } */
|
||||
28
gcc/testsuite/c-c++-common/gomp/dispatch-2.c
Normal file
28
gcc/testsuite/c-c++-common/gomp/dispatch-2.c
Normal file
@ -0,0 +1,28 @@
|
||||
/* { dg-do compile } */
|
||||
/* { dg-additional-options "-fdump-tree-gimple" } */
|
||||
|
||||
int f0 (void);
|
||||
int f1 (void);
|
||||
#pragma omp declare variant (f0) match (construct={dispatch})
|
||||
#pragma omp declare variant (f1) match (implementation={vendor(gnu)})
|
||||
int f2 (void);
|
||||
|
||||
int test (void)
|
||||
{
|
||||
int a;
|
||||
#pragma omp dispatch
|
||||
a = f2 ();
|
||||
#pragma omp dispatch novariants(1)
|
||||
a = f2 ();
|
||||
#pragma omp dispatch novariants(0)
|
||||
a = f2 ();
|
||||
#pragma omp dispatch nocontext(1)
|
||||
a = f2 ();
|
||||
#pragma omp dispatch nocontext(0)
|
||||
a = f2 ();
|
||||
return a;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 3 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } } */
|
||||
12
gcc/testsuite/c-c++-common/gomp/dispatch-3.c
Normal file
12
gcc/testsuite/c-c++-common/gomp/dispatch-3.c
Normal file
@ -0,0 +1,12 @@
|
||||
/* { dg-do compile } */
|
||||
/* { dg-additional-options "-fdump-tree-gimple" } */
|
||||
|
||||
void f2 (void);
|
||||
|
||||
void test (void)
|
||||
{
|
||||
#pragma omp dispatch /* { dg-final { scan-tree-dump-not "#pragma omp task" "gimple" } } */
|
||||
f2 ();
|
||||
#pragma omp dispatch nowait /* { dg-final { scan-tree-dump-not "nowait" "gimple" } } */
|
||||
f2 ();
|
||||
}
|
||||
18
gcc/testsuite/c-c++-common/gomp/dispatch-4.c
Normal file
18
gcc/testsuite/c-c++-common/gomp/dispatch-4.c
Normal file
@ -0,0 +1,18 @@
|
||||
/* { dg-do compile } */
|
||||
/* { dg-additional-options "-fdump-tree-gimple" } */
|
||||
|
||||
void f2 (int a);
|
||||
|
||||
void test (void)
|
||||
{
|
||||
int a;
|
||||
|
||||
#pragma omp dispatch device(-25373654)
|
||||
/* { dg-final { scan-tree-dump-times "__builtin_omp_set_default_device \\(-25373654\\);" 1 "gimple" } } */
|
||||
f2 (a);
|
||||
#pragma omp dispatch device(a + a)
|
||||
/* { dg-final { scan-tree-dump-times "(D\.\[0-9]+) = a \\* 2;.*#pragma omp dispatch.*__builtin_omp_set_default_device \\(\\1\\);.*f2 \\(a\\)" 2 "gimple" } } */
|
||||
f2 (a);
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump-times "(D\.\[0-9]+) = __builtin_omp_get_default_device \\(\\);.*__builtin_omp_set_default_device \\(\\1\\);" 4 "gimple" } } */
|
||||
34
gcc/testsuite/c-c++-common/gomp/dispatch-5.c
Normal file
34
gcc/testsuite/c-c++-common/gomp/dispatch-5.c
Normal file
@ -0,0 +1,34 @@
|
||||
/* { dg-do compile } */
|
||||
/* { dg-additional-options "-fdump-tree-gimple" } */
|
||||
|
||||
void f1 (void* p, int arr[]);
|
||||
#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: p, arr)
|
||||
void f2 (void* p, int arr[]);
|
||||
|
||||
void test (void)
|
||||
{
|
||||
void *p;
|
||||
int arr[2];
|
||||
|
||||
/* Note there are multiple matches because every variable capturing matches in addition,
|
||||
i.e. scan-tree-dump-times = 1 plus number of captures used for backward references.
|
||||
|
||||
For the first scan-tree-dump, on some targets the __builtin_omp_get_mapped_ptr get
|
||||
swapped. As the counting does not work well with variable-name capturing, it only
|
||||
uses 'scan-tree-dump' and not 'scan-tree-dump-times'. */
|
||||
|
||||
#pragma omp dispatch
|
||||
/* { dg-final { scan-tree-dump "#pragma omp dispatch\[ \t\n\r\{]*int (D\.\[0-9]+);\[ \t\n\r]*void \\* (D\.\[0-9]+);\[ \t\n\r]*void \\* (D\.\[0-9]+);\[ \t\n\r]*\\1 = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*\\2 = __builtin_omp_get_mapped_ptr \\(&arr, \\1\\);\[ \t\n\r]*\\3 = __builtin_omp_get_mapped_ptr \\(p, \\1\\);\[ \t\n\r]*f1 \\(\\3, \\2\\);|#pragma omp dispatch\[ \t\n\r\{]*int (D\.\[0-9]+);\[ \t\n\r]*void \\* (D\.\[0-9]+);\[ \t\n\r]*void \\* (D\.\[0-9]+);\[ \t\n\r]*\\4 = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*\\5 = __builtin_omp_get_mapped_ptr \\(p, \\4\\);\[ \t\n\r]*\\6 = __builtin_omp_get_mapped_ptr \\(&arr, \\4\\);\[ \t\n\r]*f1 \\(\\5, \\6\\);" "gimple" } } */
|
||||
f2 (p, arr);
|
||||
#pragma omp dispatch is_device_ptr(p)
|
||||
/* { dg-final { scan-tree-dump-times "#pragma omp dispatch is_device_ptr\\(p\\)\[ \t\n\r\{]*int (D\.\[0-9]+);\[ \t\n\r]*void \\* (D\.\[0-9]+);\[ \t\n\r]*\\1 = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*\\2 = __builtin_omp_get_mapped_ptr \\(&arr, \\1\\);\[ \t\n\r]*f1 \\(p, \\2\\);" 3 "gimple" } } */
|
||||
f2 (p, arr);
|
||||
#pragma omp dispatch is_device_ptr(arr)
|
||||
/* { dg-final { scan-tree-dump-times "#pragma omp dispatch is_device_ptr\\(arr\\)\[ \t\n\r\{]*int (D\.\[0-9]+);\[ \t\n\r]*void \\* (D\.\[0-9]+);\[ \t\n\r]*\\1 = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*\\2 = __builtin_omp_get_mapped_ptr \\(p, \\1\\);\[ \t\n\r]*f1 \\(\\2, &arr\\);" 3 "gimple" } } */
|
||||
f2 (p, arr);
|
||||
#pragma omp dispatch is_device_ptr(p, arr)
|
||||
/* { dg-final { scan-tree-dump-times "#pragma omp dispatch is_device_ptr\\(arr\\) is_device_ptr\\(p\\)\[ \t\n\r\{]*f1 \\(p, &arr\\);" 1 "gimple" } } */
|
||||
f2 (p, arr);
|
||||
}
|
||||
|
||||
|
||||
18
gcc/testsuite/c-c++-common/gomp/dispatch-6.c
Normal file
18
gcc/testsuite/c-c++-common/gomp/dispatch-6.c
Normal file
@ -0,0 +1,18 @@
|
||||
/* { dg-do compile } */
|
||||
/* { dg-additional-options "-fdump-tree-ompexp" } */
|
||||
|
||||
void f2 (void* p);
|
||||
|
||||
void test (void)
|
||||
{
|
||||
void *p;
|
||||
|
||||
#pragma omp dispatch
|
||||
/* { dg-final { scan-tree-dump-not "__builtin_GOMP_task " "ompexp" } } */
|
||||
f2 (p);
|
||||
#pragma omp dispatch depend(inout: p)
|
||||
/* { dg-final { scan-tree-dump-times "(D\.\[0-9]+)\\\[2] = &p;\[ \n]*__builtin_GOMP_taskwait_depend \\(&\\1\\);" 2 "ompexp" } } */
|
||||
f2 (p);
|
||||
}
|
||||
|
||||
|
||||
21
gcc/testsuite/c-c++-common/gomp/dispatch-7.c
Normal file
21
gcc/testsuite/c-c++-common/gomp/dispatch-7.c
Normal file
@ -0,0 +1,21 @@
|
||||
/* { dg-do compile } */
|
||||
/* { dg-additional-options "-fdump-tree-gimple" } */
|
||||
|
||||
int f0 (void);
|
||||
int f1 (void);
|
||||
#pragma omp declare variant (f0) match (construct={dispatch})
|
||||
#pragma omp declare variant (f1) match (implementation={vendor(gnu)})
|
||||
int f2 (void);
|
||||
|
||||
int test (void)
|
||||
{
|
||||
int a, n;
|
||||
#pragma omp dispatch novariants(n < 1024) nocontext(n > 1024)
|
||||
a = f2 ();
|
||||
return a;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump-times "#pragma omp dispatch nocontext\\(0\\) novariants\\(0\\)" 1 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 1 "gimple" } } */
|
||||
63
gcc/testsuite/c-c++-common/gomp/dispatch-8.c
Normal file
63
gcc/testsuite/c-c++-common/gomp/dispatch-8.c
Normal file
@ -0,0 +1,63 @@
|
||||
int g1() { return 100; }
|
||||
int g2() { return 200; }
|
||||
|
||||
#pragma omp declare variant (g1) match (construct={dispatch},user={condition(1)})
|
||||
#pragma omp declare variant (g2) match (user={condition(1)})
|
||||
int g() { return 300; }
|
||||
|
||||
|
||||
int f1(int x) { return 1 + x; }
|
||||
int f2(int x) { return 2 + x; }
|
||||
|
||||
#pragma omp declare variant (f1) match (construct={dispatch},user={condition(1)})
|
||||
#pragma omp declare variant (f2) match (user={condition(1)})
|
||||
int f(int x) { return 3 + x; }
|
||||
|
||||
|
||||
void
|
||||
test ()
|
||||
{
|
||||
int res;
|
||||
|
||||
// Call f2(g2()) due to user condition(1)
|
||||
res = f (g ());
|
||||
__builtin_printf("%d\n", res);
|
||||
if (res != 202)
|
||||
__builtin_abort ();
|
||||
|
||||
// Call 'f1(g1())' due to construct 'dispatch' and user conditional(1)
|
||||
#pragma omp dispatch
|
||||
res = f (g ());
|
||||
__builtin_printf("%d\n", res);
|
||||
if (res != 101)
|
||||
__builtin_abort ();
|
||||
|
||||
// Call 'f2(g2())' due to nocontext (i.e. ignore construct 'dispatch') and user conditional(1)
|
||||
#pragma omp dispatch nocontext(1)
|
||||
res = f (g ());
|
||||
__builtin_printf("%d\n", res);
|
||||
if (res != 202)
|
||||
__builtin_abort ();
|
||||
|
||||
// Call 'f' due to 'novariants'
|
||||
// Call 'g1' due to construct 'dispatch' and user conditional(1)
|
||||
#pragma omp dispatch novariants(1)
|
||||
res = f (g ());
|
||||
__builtin_printf("%d\n", res); /* ACTUAL RESULT: 303, i.e. 'g' and not 'g1' was called */
|
||||
if (res != 103)
|
||||
__builtin_abort ();
|
||||
|
||||
// Call 'f' due to 'novariants'
|
||||
// Call 'g2' due to nocontext (i.e. ignore construct 'dispatch') and user conditional(1)
|
||||
#pragma omp dispatch novariants(1) nocontext(1)
|
||||
res = f (g ());
|
||||
__builtin_printf("%d\n", res); /* ACTUAL RESULT: 303, i.e. 'g' and not 'g2' was called */
|
||||
if (res != 203)
|
||||
__builtin_abort ();
|
||||
}
|
||||
|
||||
int
|
||||
main ()
|
||||
{
|
||||
test ();
|
||||
}
|
||||
17
gcc/testsuite/c-c++-common/gomp/dispatch-9.c
Normal file
17
gcc/testsuite/c-c++-common/gomp/dispatch-9.c
Normal file
@ -0,0 +1,17 @@
|
||||
/* { dg-do compile } */
|
||||
/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */
|
||||
|
||||
/* Check that the right call to f is wrapped in a GOMP_DISPATCH internal function
|
||||
before translation and that it is stripped during gimplification. */
|
||||
|
||||
int f(int);
|
||||
void g(int *x)
|
||||
{
|
||||
#pragma omp dispatch
|
||||
x[f(1)] = f(f(2));
|
||||
// ^ only this call to f is a dispatch call
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump "\.GOMP_DISPATCH \\(f \\(f \\(2\\)\\)\\)" "original" } } */
|
||||
/* { dg-final { scan-tree-dump-times "\.GOMP_DISPATCH" 1 "original" } } */
|
||||
/* { dg-final { scan-tree-dump-not "\.GOMP_DISPATCH" "gimple" } } */
|
||||
76
libgomp/testsuite/libgomp.c-c++-common/dispatch-1.c
Normal file
76
libgomp/testsuite/libgomp.c-c++-common/dispatch-1.c
Normal file
@ -0,0 +1,76 @@
|
||||
// Adapted from OpenMP examples
|
||||
|
||||
#include <omp.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
int baz (double *d_bv, const double *d_av, int n)
|
||||
{
|
||||
#pragma omp distribute parallel for
|
||||
for (int i = 0; i < n; i++)
|
||||
d_bv[i] = d_av[i] * i;
|
||||
return -3;
|
||||
}
|
||||
|
||||
int bar (double *d_bv, const double *d_av, int n)
|
||||
{
|
||||
#pragma omp target is_device_ptr(d_bv, d_av)
|
||||
for (int i = 0; i < n; i++)
|
||||
d_bv[i] = d_av[i] * i;
|
||||
return -2;
|
||||
}
|
||||
|
||||
#pragma omp declare variant(bar) match(construct={dispatch}) adjust_args(need_device_ptr: bv, av)
|
||||
#pragma omp declare variant(baz) match(implementation={vendor(gnu)})
|
||||
int foo (double *bv, const double *av, int n)
|
||||
{
|
||||
for (int i = 0; i < n; i++)
|
||||
bv[i] = av[i] * i;
|
||||
return -1;
|
||||
}
|
||||
|
||||
int test (int n)
|
||||
{
|
||||
const double e = 2.71828;
|
||||
|
||||
double *av = (double *) malloc (n * sizeof (*av));
|
||||
double *bv = (double *) malloc (n * sizeof (*bv));
|
||||
double *d_bv = (double *) malloc (n * sizeof (*d_bv));
|
||||
|
||||
for (int i = 0; i < n; i++)
|
||||
{
|
||||
av[i] = e * i;
|
||||
bv[i] = 0.0;
|
||||
d_bv[i] = 0.0;
|
||||
}
|
||||
|
||||
int f, last_dev = omp_get_num_devices () - 1;
|
||||
#pragma omp target data map(to: av[:n]) map(from: d_bv[:n]) device(last_dev) if (n == 1024)
|
||||
{
|
||||
#pragma omp dispatch nocontext(n > 1024) novariants(n < 1024) device(last_dev)
|
||||
f = foo (d_bv, av, n);
|
||||
}
|
||||
|
||||
foo (bv, av, n);
|
||||
for (int i = 0; i < n; i++)
|
||||
{
|
||||
if (d_bv[i] != bv[i])
|
||||
{
|
||||
fprintf (stderr, "ERROR at %d: %lf (act) != %lf (exp)\n", i, d_bv[i], bv[i]);
|
||||
return 1;
|
||||
}
|
||||
}
|
||||
return f;
|
||||
}
|
||||
|
||||
int
|
||||
main (void)
|
||||
{
|
||||
int ret = test(1023);
|
||||
if (ret != -1) return 1;
|
||||
ret = test(1024);
|
||||
if (ret != -2) return 1;
|
||||
ret = test(1025);
|
||||
if (ret != -3) return 1;
|
||||
return 0;
|
||||
}
|
||||
84
libgomp/testsuite/libgomp.c-c++-common/dispatch-2.c
Normal file
84
libgomp/testsuite/libgomp.c-c++-common/dispatch-2.c
Normal file
@ -0,0 +1,84 @@
|
||||
// Adapted from OpenMP examples
|
||||
|
||||
#include <omp.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
int baz (double *d_bv, const double *d_av, int n);
|
||||
int bar (double *d_bv, const double *d_av, int n);
|
||||
|
||||
#pragma omp declare variant(bar) match(construct={dispatch}) adjust_args(need_device_ptr: f_bv, f_av)
|
||||
#pragma omp declare variant(baz) match(implementation={vendor(gnu)})
|
||||
int foo (double *f_bv, const double *f_av, int n);
|
||||
|
||||
int baz (double *bv, const double *av, int n);
|
||||
int bar (double *bv, const double *av, int n);
|
||||
|
||||
int foo (double *bv, const double *av, int n)
|
||||
{
|
||||
for (int i = 0; i < n; i++)
|
||||
bv[i] = av[i] * i;
|
||||
return -1;
|
||||
}
|
||||
|
||||
int baz (double *d_bv, const double *d_av, int n)
|
||||
{
|
||||
#pragma omp distribute parallel for
|
||||
for (int i = 0; i < n; i++)
|
||||
d_bv[i] = d_av[i] * i;
|
||||
return -3;
|
||||
}
|
||||
|
||||
int bar (double *d_bv, const double *d_av, int n)
|
||||
{
|
||||
#pragma omp target is_device_ptr(d_bv, d_av)
|
||||
for (int i = 0; i < n; i++)
|
||||
d_bv[i] = d_av[i] * i;
|
||||
return -2;
|
||||
}
|
||||
|
||||
int test (int n)
|
||||
{
|
||||
const double e = 2.71828;
|
||||
|
||||
double *av = (double *) malloc (n * sizeof (*av));
|
||||
double *bv = (double *) malloc (n * sizeof (*bv));
|
||||
double *d_bv = (double *) malloc (n * sizeof (*d_bv));
|
||||
|
||||
for (int i = 0; i < n; i++)
|
||||
{
|
||||
av[i] = e * i;
|
||||
bv[i] = 0.0;
|
||||
d_bv[i] = 0.0;
|
||||
}
|
||||
|
||||
int f, last_dev = omp_get_num_devices () - 1;
|
||||
#pragma omp target data map(to: av[:n]) map(from: d_bv[:n]) device(last_dev) if (n == 1024)
|
||||
{
|
||||
#pragma omp dispatch nocontext(n > 1024) novariants(n < 1024) device(last_dev)
|
||||
f = foo (d_bv, av, n);
|
||||
}
|
||||
|
||||
foo (bv, av, n);
|
||||
for (int i = 0; i < n; i++)
|
||||
{
|
||||
if (d_bv[i] != bv[i])
|
||||
{
|
||||
fprintf (stderr, "ERROR at %d: %lf (act) != %lf (exp)\n", i, d_bv[i], bv[i]);
|
||||
return 1;
|
||||
}
|
||||
}
|
||||
return f;
|
||||
}
|
||||
|
||||
int
|
||||
main (void)
|
||||
{
|
||||
int ret = test(1023);
|
||||
if (ret != -1) return 1;
|
||||
ret = test(1024);
|
||||
if (ret != -2) return 1;
|
||||
ret = test(1025);
|
||||
if (ret != -3) return 1;
|
||||
return 0;
|
||||
}
|
||||
Loading…
Reference in New Issue
Block a user