diff options
Diffstat (limited to 'gcc/testsuite/c-c++-common/gomp')
-rw-r--r-- | gcc/testsuite/c-c++-common/gomp/adjust-args-3.c | 85 |
1 files changed, 85 insertions, 0 deletions
diff --git a/gcc/testsuite/c-c++-common/gomp/adjust-args-3.c b/gcc/testsuite/c-c++-common/gomp/adjust-args-3.c new file mode 100644 index 0000000..f62272c --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/adjust-args-3.c @@ -0,0 +1,85 @@ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +// Do diagnostic check / dump check only; +// Note: this test should work as run-test as well. + +#if 0 + #include <omp.h> +#else + #ifdef __cplusplus + extern "C" { + #endif + extern int omp_get_default_device (); + extern int omp_get_num_devices (); + #ifdef __cplusplus + } + #endif +#endif + + +void f(int *x, int *y); +#pragma omp declare variant(f) adjust_args(need_device_ptr: x, y) match(construct={dispatch}) +void g(int *x, int *y); + +void +sub (int *a, int *b) +{ + // The has_device_addr is a bit questionable as the caller is not actually + // passing a device address - but we cannot pass one because of the + // following: + // + // As for 'b' need_device_ptr has been specified and 'b' is not + // in the semantic requirement set 'is_device_ptr' (and only in 'has_device_addr') + // "the argument is converted in the same manner that a use_device_ptr clause + // on a target_data construct converts its pointer" + #pragma omp dispatch is_device_ptr(a), has_device_addr(b) /* { dg-warning "'has_device_addr' for 'b' does not imply 'is_device_ptr' required for 'need_device_ptr' \\\[-Wopenmp\\\]" } */ + g(a, b); +} + +void +f(int *from, int *to) +{ + static int cnt = 0; + cnt++; + if (cnt >= 3) + { + if (omp_get_default_device () != -1 + && omp_get_default_device () < omp_get_num_devices ()) + { + // On offload device but not mapped + if (from != (void *)0L) // Not mapped + __builtin_abort (); + } + else if (from[0] != 5) + __builtin_abort (); + return; + } + #pragma omp target is_device_ptr(from, to) + { + to[0] = from[0] * 10; + to[1] = from[1] * 10; + } +} + +int +main () +{ + int A[2], B[2] = {123, 456}, C[1] = {5}; + int *p = A; + #pragma omp target enter data map(A, B) + + /* Note: We don't add 'use_device_addr(B)' here; + if we do, it will fail with an illegal memory access (why?). */ + #pragma omp target data use_device_ptr(p) + { + sub(p, B); + sub(C, B); /* C is not mapped -> 'from' ptr == NULL */ + } + + #pragma omp target exit data map(A, B) +} + +// { dg-final { scan-tree-dump-times "#pragma omp dispatch has_device_addr\\(b\\) is_device_ptr\\(a\\)" 1 "gimple" } } +// { dg-final { scan-tree-dump-times "__builtin_omp_get_mapped_ptr" 1 "gimple" } } +// { dg-final { scan-tree-dump-times "D\\.\[0-9\]+ = __builtin_omp_get_mapped_ptr \\(b" 1 "gimple" } } +// { dg-final { scan-tree-dump-times "f \\(a, D\\.\[0-9\]+\\);" 1 "gimple" } } |