diff options
author | Paul-Antoine Arras <parras@baylibre.com> | 2024-11-20 15:28:58 +0100 |
---|---|---|
committer | Paul-Antoine Arras <parras@baylibre.com> | 2024-11-20 15:31:22 +0100 |
commit | 377eff7c38e8df1388b0a1eeb6afdbcaf12c3551 (patch) | |
tree | 9790d2e16ad6537f2eaebfcc2e900c3630d018c0 /libgomp/testsuite/libgomp.c-c++-common/dispatch-2.c | |
parent | ed49709acda4938f71ef2dfac68450be5429db3d (diff) | |
download | gcc-377eff7c38e8df1388b0a1eeb6afdbcaf12c3551.zip gcc-377eff7c38e8df1388b0a1eeb6afdbcaf12c3551.tar.gz gcc-377eff7c38e8df1388b0a1eeb6afdbcaf12c3551.tar.bz2 |
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.
Diffstat (limited to 'libgomp/testsuite/libgomp.c-c++-common/dispatch-2.c')
-rw-r--r-- | libgomp/testsuite/libgomp.c-c++-common/dispatch-2.c | 84 |
1 files changed, 84 insertions, 0 deletions
diff --git a/libgomp/testsuite/libgomp.c-c++-common/dispatch-2.c b/libgomp/testsuite/libgomp.c-c++-common/dispatch-2.c new file mode 100644 index 0000000..faa0d8a --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/dispatch-2.c @@ -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; +} |