diff options
Diffstat (limited to 'libgomp/testsuite/libgomp.c')
49 files changed, 3265 insertions, 33 deletions
diff --git a/libgomp/testsuite/libgomp.c/affinity-2.c b/libgomp/testsuite/libgomp.c/affinity-2.c new file mode 100644 index 0000000..f821657 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/affinity-2.c @@ -0,0 +1,89 @@ +/* { dg-do run } */ +/* { dg-set-target-env-var OMP_PROC_BIND "spread,close" } */ +/* { dg-set-target-env-var OMP_PLACES "{6,7}:4:-2,!{2,3}" } */ +/* { dg-set-target-env-var OMP_NUM_THREADS "2" } */ + +#include <omp.h> +#include <stdlib.h> +#include <stdio.h> + +int * +get_buf (int nump) +{ + static int *buf; + static size_t buf_size; + if ((size_t) nump > buf_size) + { + buf_size *= 2; + if (nump > buf_size) + buf_size = nump + 64; + int *bufn = realloc (buf, buf_size * sizeof (int)); + if (bufn == NULL) + { + fprintf (stderr, "memory allocation error\n"); + exit (1); + } + buf = bufn; + } + return buf; +} + +void +print_place (int count, int *ids) +{ + int i, j; + printf ("{"); + for (i = 0; i < count; i++) + { + for (j = i + 1; j < count; j++) + if (ids[j] != ids[i] + (j - i)) + break; + if (i) + printf (","); + if (j == i + 1) + printf ("%d", ids[i]); + else + { + printf ("%d:%d", ids[i], j - i); + i = j - 1; + } + } + printf ("}\n"); +} + +void +print_place_var (void) +{ + int place = omp_get_place_num (); + int num_places = omp_get_partition_num_places (); + int *ids = get_buf (num_places); + omp_get_partition_place_nums (ids); + printf ("place %d\n", place); + if (num_places) + printf ("partition %d-%d\n", ids[0], ids[num_places - 1]); +} + +int +main () +{ + int i, num = omp_get_num_places (), nump, *ids; + printf ("omp_get_num_places () == %d\n", num); + for (i = 0; i < num; i++) + { + printf ("place %d ", i); + nump = omp_get_place_num_procs (i); + ids = get_buf (nump); + omp_get_place_proc_ids (i, ids); + print_place (nump, ids); + } + print_place_var (); + omp_set_nested (1); + #pragma omp parallel + if (omp_get_thread_num () == omp_get_num_threads () - 1) + { + #pragma omp parallel + if (omp_get_thread_num () == omp_get_num_threads () - 1) + print_place_var (); + } + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/doacross-1.c b/libgomp/testsuite/libgomp.c/doacross-1.c new file mode 100644 index 0000000..0794c80 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/doacross-1.c @@ -0,0 +1,181 @@ +extern void abort (void); + +#define N 256 +int a[N], b[N / 16][8][4], c[N / 32][8][8]; +volatile int d, e; + +int +main () +{ + int i, j, k, l, m; + #pragma omp parallel private (l) + { + #pragma omp for schedule(static, 1) ordered (1) nowait + for (i = 0; i < N; i++) + { + #pragma omp atomic write + a[i] = 1; + #pragma omp ordered depend(sink: i - 1) + if (i) + { + #pragma omp atomic read + l = a[i - 1]; + if (l < 2) + abort (); + } + #pragma omp atomic write + a[i] = 2; + if (i < N - 1) + { + #pragma omp atomic read + l = a[i + 1]; + if (l == 3) + abort (); + } + #pragma omp ordered depend(source) + #pragma omp atomic write + a[i] = 3; + } + #pragma omp for schedule(static, 0) ordered (3) nowait + for (i = 2; i < N / 16 - 1; i++) + for (j = 0; j < 8; j += 2) + for (k = 1; k <= 3; k++) + { + #pragma omp atomic write + b[i][j][k] = 1; + #pragma omp ordered depend(sink: i, j - 2, k - 1) \ + depend(sink: i - 2, j - 2, k + 1) + #pragma omp ordered depend(sink: i - 3, j + 2, k - 2) + if (j >= 2 && k > 1) + { + #pragma omp atomic read + l = b[i][j - 2][k - 1]; + if (l < 2) + abort (); + } + #pragma omp atomic write + b[i][j][k] = 2; + if (i >= 4 && j >= 2 && k < 3) + { + #pragma omp atomic read + l = b[i - 2][j - 2][k + 1]; + if (l < 2) + abort (); + } + if (i >= 5 && j < N / 16 - 3 && k == 3) + { + #pragma omp atomic read + l = b[i - 3][j + 2][k - 2]; + if (l < 2) + abort (); + } + #pragma omp ordered depend(source) + #pragma omp atomic write + b[i][j][k] = 3; + } +#define A(n) int n; +#define B(n) A(n##0) A(n##1) A(n##2) A(n##3) +#define C(n) B(n##0) B(n##1) B(n##2) B(n##3) +#define D(n) C(n##0) C(n##1) C(n##2) C(n##3) + D(m) +#undef A + #pragma omp for collapse (2) ordered(61) schedule(dynamic, 15) + for (i = 0; i < N / 32; i++) + for (j = 7; j > 1; j--) + for (k = 6; k >= 0; k -= 2) +#define A(n) for (n = 4; n < 5; n++) + D(m) +#undef A + { + #pragma omp atomic write + c[i][j][k] = 1; +#define A(n) ,n +#define E(n) C(n##0) C(n##1) C(n##2) B(n##30) B(n##31) A(n##320) A(n##321) + #pragma omp ordered depend (sink: i, j, k + 2 E(m)) \ + depend (sink:i - 2, j + 1, k - 4 E(m)) \ + depend(sink: i - 1, j - 2, k - 2 E(m)) + if (k <= 4) + { + l = c[i][j][k + 2]; + if (l < 2) + abort (); + } + #pragma omp atomic write + c[i][j][k] = 2; + if (i >= 2 && j < 7 && k >= 4) + { + l = c[i - 2][j + 1][k - 4]; + if (l < 2) + abort (); + } + if (i >= 1 && j >= 4 && k >= 2) + { + l = c[i - 1][j - 2][k - 2]; + if (l < 2) + abort (); + } + #pragma omp ordered depend (source) + #pragma omp atomic write + c[i][j][k] = 3; + } + + #pragma omp for collapse(2) ordered(4) lastprivate (i, j, k) + for (i = 0; i < d + 1; i++) + for (j = d + 1; j >= 0; j--) + for (k = 0; k < d; k++) + for (l = 0; l < d + 2; l++) + { + #pragma omp ordered depend (source) + #pragma omp ordered depend (sink:i - 2, j + 2, k - 2, l) + if (!e) + abort (); + } + #pragma omp single + { + if (i != 1 || j != -1 || k != 0) + abort (); + i = 8; j = 9; k = 10; + } + #pragma omp for collapse(2) ordered(4) lastprivate (i, j, k, m) + for (i = 0; i < d + 1; i++) + for (j = d + 1; j >= 0; j--) + for (k = 0; k < d + 2; k++) + for (m = 0; m < d; m++) + { + #pragma omp ordered depend (source) + #pragma omp ordered depend (sink:i - 2, j + 2, k - 2, m) + abort (); + } + #pragma omp single + if (i != 1 || j != -1 || k != 2 || m != 0) + abort (); + #pragma omp for collapse(2) ordered(4) nowait + for (i = 0; i < d + 1; i++) + for (j = d; j > 0; j--) + for (k = 0; k < d + 2; k++) + for (l = 0; l < d + 4; l++) + { + #pragma omp ordered depend (source) + #pragma omp ordered depend (sink:i - 2, j + 2, k - 2, l) + if (!e) + abort (); + } + #pragma omp for nowait + for (i = 0; i < N; i++) + if (a[i] != 3) + abort (); + #pragma omp for collapse(2) private(k) nowait + for (i = 0; i < N / 16; i++) + for (j = 0; j < 8; j++) + for (k = 0; k < 4; k++) + if (b[i][j][k] != 3 * (i >= 2 && i < N / 16 - 1 && (j & 1) == 0 && k >= 1)) + abort (); + #pragma omp for collapse(3) nowait + for (i = 0; i < N / 32; i++) + for (j = 0; j < 8; j++) + for (k = 0; k < 8; k++) + if (c[i][j][k] != 3 * (j >= 2 && (k & 1) == 0)) + abort (); + } + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/doacross-2.c b/libgomp/testsuite/libgomp.c/doacross-2.c new file mode 100644 index 0000000..e491bb2 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/doacross-2.c @@ -0,0 +1,225 @@ +extern void abort (void); + +#define N 256 +int a[N], b[N / 16][8][4], c[N / 32][8][8], g[N / 16][8][6]; +volatile int d, e; +volatile unsigned long long f; + +int +main () +{ + unsigned long long i; + int j, k, l, m; + #pragma omp parallel private (l) + { + #pragma omp for schedule(static, 1) ordered (1) nowait + for (i = 1; i < N + f; i++) + { + #pragma omp atomic write + a[i] = 1; + #pragma omp ordered depend(sink: i - 1) + if (i > 1) + { + #pragma omp atomic read + l = a[i - 1]; + if (l < 2) + abort (); + } + #pragma omp atomic write + a[i] = 2; + if (i < N - 1) + { + #pragma omp atomic read + l = a[i + 1]; + if (l == 3) + abort (); + } + #pragma omp ordered depend(source) + #pragma omp atomic write + a[i] = 3; + } + #pragma omp for schedule(static, 0) ordered (3) nowait + for (i = 3; i < N / 16 - 1 + f; i++) + for (j = 0; j < 8; j += 2) + for (k = 1; k <= 3; k++) + { + #pragma omp atomic write + b[i][j][k] = 1; + #pragma omp ordered depend(sink: i, j - 2, k - 1) \ + depend(sink: i - 2, j - 2, k + 1) + #pragma omp ordered depend(sink: i - 3, j + 2, k - 2) + if (j >= 2 && k > 1) + { + #pragma omp atomic read + l = b[i][j - 2][k - 1]; + if (l < 2) + abort (); + } + #pragma omp atomic write + b[i][j][k] = 2; + if (i >= 5 && j >= 2 && k < 3) + { + #pragma omp atomic read + l = b[i - 2][j - 2][k + 1]; + if (l < 2) + abort (); + } + if (i >= 6 && j < N / 16 - 3 && k == 3) + { + #pragma omp atomic read + l = b[i - 3][j + 2][k - 2]; + if (l < 2) + abort (); + } + #pragma omp ordered depend(source) + #pragma omp atomic write + b[i][j][k] = 3; + } +#define A(n) int n; +#define B(n) A(n##0) A(n##1) A(n##2) A(n##3) +#define C(n) B(n##0) B(n##1) B(n##2) B(n##3) +#define D(n) C(n##0) C(n##1) C(n##2) C(n##3) + D(m) +#undef A + #pragma omp for collapse (2) ordered(61) schedule(dynamic, 15) + for (i = 2; i < N / 32 + f; i++) + for (j = 7; j > 1; j--) + for (k = 6; k >= 0; k -= 2) +#define A(n) for (n = 4; n < 5; n++) + D(m) +#undef A + { + #pragma omp atomic write + c[i][j][k] = 1; +#define A(n) ,n +#define E(n) C(n##0) C(n##1) C(n##2) B(n##30) B(n##31) A(n##320) A(n##321) + #pragma omp ordered depend (sink: i, j, k + 2 E(m)) \ + depend (sink:i - 2, j + 1, k - 4 E(m)) \ + depend(sink: i - 1, j - 2, k - 2 E(m)) + if (k <= 4) + { + l = c[i][j][k + 2]; + if (l < 2) + abort (); + } + #pragma omp atomic write + c[i][j][k] = 2; + if (i >= 4 && j < 7 && k >= 4) + { + l = c[i - 2][j + 1][k - 4]; + if (l < 2) + abort (); + } + if (i >= 3 && j >= 4 && k >= 2) + { + l = c[i - 1][j - 2][k - 2]; + if (l < 2) + abort (); + } + #pragma omp ordered depend (source) + #pragma omp atomic write + c[i][j][k] = 3; + } + #pragma omp for schedule(static, 0) ordered (3) nowait + for (j = 0; j < N / 16 - 1; j++) + for (k = 0; k < 8; k += 2) + for (i = 3; i <= 5 + f; i++) + { + #pragma omp atomic write + g[j][k][i] = 1; + #pragma omp ordered depend(sink: j, k - 2, i - 1) \ + depend(sink: j - 2, k - 2, i + 1) + #pragma omp ordered depend(sink: j - 3, k + 2, i - 2) + if (k >= 2 && i > 3) + { + #pragma omp atomic read + l = g[j][k - 2][i - 1]; + if (l < 2) + abort (); + } + #pragma omp atomic write + g[j][k][i] = 2; + if (j >= 2 && k >= 2 && i < 5) + { + #pragma omp atomic read + l = g[j - 2][k - 2][i + 1]; + if (l < 2) + abort (); + } + if (j >= 3 && k < N / 16 - 3 && i == 5) + { + #pragma omp atomic read + l = g[j - 3][k + 2][i - 2]; + if (l < 2) + abort (); + } + #pragma omp ordered depend(source) + #pragma omp atomic write + g[j][k][i] = 3; + } + #pragma omp for collapse(2) ordered(4) lastprivate (i, j, k) + for (i = 2; i < f + 3; i++) + for (j = d + 1; j >= 0; j--) + for (k = 0; k < d; k++) + for (l = 0; l < d + 2; l++) + { + #pragma omp ordered depend (source) + #pragma omp ordered depend (sink:i - 2, j + 2, k - 2, l) + if (!e) + abort (); + } + #pragma omp single + { + if (i != 3 || j != -1 || k != 0) + abort (); + i = 8; j = 9; k = 10; + } + #pragma omp for collapse(2) ordered(4) lastprivate (i, j, k, m) + for (i = 2; i < f + 3; i++) + for (j = d + 1; j >= 0; j--) + for (k = 0; k < d + 2; k++) + for (m = 0; m < d; m++) + { + #pragma omp ordered depend (source) + #pragma omp ordered depend (sink:i - 2, j + 2, k - 2, m) + abort (); + } + #pragma omp single + if (i != 3 || j != -1 || k != 2 || m != 0) + abort (); + #pragma omp for collapse(2) ordered(4) nowait + for (i = 2; i < f + 3; i++) + for (j = d; j > 0; j--) + for (k = 0; k < d + 2; k++) + for (l = 0; l < d + 4; l++) + { + #pragma omp ordered depend (source) + #pragma omp ordered depend (sink:i - 2, j + 2, k - 2, l) + if (!e) + abort (); + } + #pragma omp for nowait + for (i = 0; i < N; i++) + if (a[i] != 3 * (i >= 1)) + abort (); + #pragma omp for collapse(2) private(k) nowait + for (i = 0; i < N / 16; i++) + for (j = 0; j < 8; j++) + for (k = 0; k < 4; k++) + if (b[i][j][k] != 3 * (i >= 3 && i < N / 16 - 1 && (j & 1) == 0 && k >= 1)) + abort (); + #pragma omp for collapse(3) nowait + for (i = 0; i < N / 32; i++) + for (j = 0; j < 8; j++) + for (k = 0; k < 8; k++) + if (c[i][j][k] != 3 * (i >= 2 && j >= 2 && (k & 1) == 0)) + abort (); + #pragma omp for collapse(2) private(k) nowait + for (i = 0; i < N / 16; i++) + for (j = 0; j < 8; j++) + for (k = 0; k < 6; k++) + if (g[i][j][k] != 3 * (i < N / 16 - 1 && (j & 1) == 0 && k >= 3)) + abort (); + } + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/examples-4/declare_target-1.c b/libgomp/testsuite/libgomp.c/examples-4/declare_target-1.c index beca855..6d4bc4f 100644 --- a/libgomp/testsuite/libgomp.c/examples-4/declare_target-1.c +++ b/libgomp/testsuite/libgomp.c/examples-4/declare_target-1.c @@ -20,7 +20,7 @@ int fib_wrapper (int n) { int x = 0; - #pragma omp target if(n > THRESHOLD) + #pragma omp target if(n > THRESHOLD) map(from:x) x = fib (n); return x; diff --git a/libgomp/testsuite/libgomp.c/examples-4/declare_target-4.c b/libgomp/testsuite/libgomp.c/examples-4/declare_target-4.c index db70460..f241436 100644 --- a/libgomp/testsuite/libgomp.c/examples-4/declare_target-4.c +++ b/libgomp/testsuite/libgomp.c/examples-4/declare_target-4.c @@ -41,7 +41,7 @@ float accum (int k) int i; float tmp = 0.0; - #pragma omp target + #pragma omp target map(tofrom:tmp) #pragma omp parallel for reduction(+:tmp) for (i = 0; i < N; i++) tmp += Pfun (i, k); diff --git a/libgomp/testsuite/libgomp.c/examples-4/declare_target-5.c b/libgomp/testsuite/libgomp.c/examples-4/declare_target-5.c index b550f1f..33d6137 100644 --- a/libgomp/testsuite/libgomp.c/examples-4/declare_target-5.c +++ b/libgomp/testsuite/libgomp.c/examples-4/declare_target-5.c @@ -48,7 +48,7 @@ float accum () int i, k; float tmp = 0.0; - #pragma omp target + #pragma omp target map(tofrom:tmp) #pragma omp parallel for reduction(+:tmp) for (i = 0; i < N; i++) { diff --git a/libgomp/testsuite/libgomp.c/examples-4/device-1.c b/libgomp/testsuite/libgomp.c/examples-4/device-1.c index f7c84fb..dad8572 100644 --- a/libgomp/testsuite/libgomp.c/examples-4/device-1.c +++ b/libgomp/testsuite/libgomp.c/examples-4/device-1.c @@ -10,11 +10,11 @@ int main () int b = 0; int c, d; - #pragma omp target if(a > 200 && a < 400) + #pragma omp target if(a > 200 && a < 400) map(from: c) c = omp_is_initial_device (); #pragma omp target data map(to: b) if(a > 200 && a < 400) - #pragma omp target + #pragma omp target map(from: b, d) { b = 100; d = omp_is_initial_device (); @@ -26,11 +26,11 @@ int main () a += 200; b = 0; - #pragma omp target if(a > 200 && a < 400) + #pragma omp target if(a > 200 && a < 400) map(from: c) c = omp_is_initial_device (); #pragma omp target data map(to: b) if(a > 200 && a < 400) - #pragma omp target + #pragma omp target map(from: b, d) { b = 100; d = omp_is_initial_device (); @@ -42,11 +42,11 @@ int main () a += 200; b = 0; - #pragma omp target if(a > 200 && a < 400) + #pragma omp target if(a > 200 && a < 400) map(from: c) c = omp_is_initial_device (); #pragma omp target data map(to: b) if(a > 200 && a < 400) - #pragma omp target + #pragma omp target map(from: b, d) { b = 100; d = omp_is_initial_device (); diff --git a/libgomp/testsuite/libgomp.c/examples-4/device-3.c b/libgomp/testsuite/libgomp.c/examples-4/device-3.c index 8a0cf7c..af086533 100644 --- a/libgomp/testsuite/libgomp.c/examples-4/device-3.c +++ b/libgomp/testsuite/libgomp.c/examples-4/device-3.c @@ -9,7 +9,7 @@ int main () int res; int default_device = omp_get_default_device (); - #pragma omp target + #pragma omp target map(from: res) res = omp_is_initial_device (); if (res) @@ -17,7 +17,7 @@ int main () omp_set_default_device (omp_get_num_devices ()); - #pragma omp target + #pragma omp target map(from: res) res = omp_is_initial_device (); if (!res) diff --git a/libgomp/testsuite/libgomp.c/examples-4/target_data-3.c b/libgomp/testsuite/libgomp.c/examples-4/target_data-3.c index abb2838..46b6740 100644 --- a/libgomp/testsuite/libgomp.c/examples-4/target_data-3.c +++ b/libgomp/testsuite/libgomp.c/examples-4/target_data-3.c @@ -47,7 +47,7 @@ void gramSchmidt (int Q[][COLS], const int rows, const int cols) { int tmp = 0; - #pragma omp target + #pragma omp target map(tofrom:tmp) #pragma omp parallel for reduction(+:tmp) for (i = 0; i < rows; i++) tmp += (Q[i][k] * Q[i][k]); diff --git a/libgomp/testsuite/libgomp.c/examples-4/teams-2.c b/libgomp/testsuite/libgomp.c/examples-4/teams-2.c index 8bbbc35..7d0a60e 100644 --- a/libgomp/testsuite/libgomp.c/examples-4/teams-2.c +++ b/libgomp/testsuite/libgomp.c/examples-4/teams-2.c @@ -32,7 +32,7 @@ float dotprod (float B[], float C[], int n, int block_size, int i, i0; float sum = 0; - #pragma omp target map(to: B[0:n], C[0:n]) + #pragma omp target map(to: B[0:n], C[0:n]) map(tofrom: sum) #pragma omp teams num_teams(num_teams) thread_limit(block_threads) \ reduction(+:sum) #pragma omp distribute diff --git a/libgomp/testsuite/libgomp.c/examples-4/teams-3.c b/libgomp/testsuite/libgomp.c/examples-4/teams-3.c index b670878..5fe63a6 100644 --- a/libgomp/testsuite/libgomp.c/examples-4/teams-3.c +++ b/libgomp/testsuite/libgomp.c/examples-4/teams-3.c @@ -31,7 +31,7 @@ float dotprod (float B[], float C[], int n) int i; float sum = 0; - #pragma omp target teams map(to: B[0:n], C[0:n]) + #pragma omp target teams map(to: B[0:n], C[0:n]) map(tofrom: sum) #pragma omp distribute parallel for reduction(+:sum) for (i = 0; i < n; i++) sum += B[i] * C[i]; diff --git a/libgomp/testsuite/libgomp.c/examples-4/teams-4.c b/libgomp/testsuite/libgomp.c/examples-4/teams-4.c index 9aef78e..6136eab 100644 --- a/libgomp/testsuite/libgomp.c/examples-4/teams-4.c +++ b/libgomp/testsuite/libgomp.c/examples-4/teams-4.c @@ -31,7 +31,7 @@ float dotprod (float B[], float C[], int n) int i; float sum = 0; - #pragma omp target map(to: B[0:n], C[0:n]) + #pragma omp target map(to: B[0:n], C[0:n]) map(tofrom:sum) #pragma omp teams num_teams(8) thread_limit(16) #pragma omp distribute parallel for reduction(+:sum) \ dist_schedule(static, 1024) \ diff --git a/libgomp/testsuite/libgomp.c/for-2.h b/libgomp/testsuite/libgomp.c/for-2.h index 920d23b..0bd116c 100644 --- a/libgomp/testsuite/libgomp.c/for-2.h +++ b/libgomp/testsuite/libgomp.c/for-2.h @@ -11,11 +11,21 @@ noreturn (void) #ifndef SC #define SC #endif +#ifndef OMPTGT +#define OMPTGT +#endif +#ifndef OMPTO +#define OMPTO(v) do {} while (0) +#endif +#ifndef OMPFROM +#define OMPFROM(v) do {} while (0) +#endif __attribute__((noinline, noclone)) void N(f0) (void) { int i; + OMPTGT #pragma omp F S for (i = 0; i < 1500; i++) a[i] += 2; @@ -24,6 +34,7 @@ N(f0) (void) __attribute__((noinline, noclone)) void N(f1) (void) { + OMPTGT #pragma omp F S for (unsigned int i = __INT_MAX__; i < 3000U + __INT_MAX__; i += 2) a[(i - __INT_MAX__) >> 1] -= 2; @@ -33,6 +44,7 @@ __attribute__((noinline, noclone)) void N(f2) (void) { unsigned long long i; + OMPTGT #pragma omp F S for (i = __LONG_LONG_MAX__ + 4500ULL - 27; i > __LONG_LONG_MAX__ - 27ULL; i -= 3) @@ -42,6 +54,7 @@ N(f2) (void) __attribute__((noinline, noclone)) void N(f3) (long long n1, long long n2, long long s3) { + OMPTGT #pragma omp F S for (long long i = n1 + 23; i > n2 - 25; i -= s3) a[i + 48] += 7; @@ -51,6 +64,7 @@ __attribute__((noinline, noclone)) void N(f4) (void) { unsigned int i; + OMPTGT #pragma omp F S for (i = 30; i < 20; i += 2) a[i] += 10; @@ -61,6 +75,7 @@ N(f5) (int n11, int n12, int n21, int n22, int n31, int n32, int s1, int s2, int s3) { SC int v1, v2, v3; + OMPTGT #pragma omp F S collapse(3) for (v1 = n11; v1 < n12; v1 += s1) for (v2 = n21; v2 < n22; v2 += s2) @@ -74,6 +89,7 @@ N(f6) (int n11, int n12, int n21, int n22, long long n31, long long n32, { SC int v1, v2; SC long long v3; + OMPTGT #pragma omp F S collapse(3) for (v1 = n11; v1 > n12; v1 += s1) for (v2 = n21; v2 > n22; v2 += s2) @@ -86,6 +102,7 @@ N(f7) (void) { SC unsigned int v1, v3; SC unsigned long long v2; + OMPTGT #pragma omp F S collapse(3) for (v1 = 0; v1 < 20; v1 += 2) for (v2 = __LONG_LONG_MAX__ + 16ULL; @@ -98,6 +115,7 @@ __attribute__((noinline, noclone)) void N(f8) (void) { SC long long v1, v2, v3; + OMPTGT #pragma omp F S collapse(3) for (v1 = 0; v1 < 20; v1 += 2) for (v2 = 30; v2 < 20; v2++) @@ -109,6 +127,7 @@ __attribute__((noinline, noclone)) void N(f9) (void) { int i; + OMPTGT #pragma omp F S for (i = 20; i < 10; i++) { @@ -122,6 +141,7 @@ __attribute__((noinline, noclone)) void N(f10) (void) { SC int i; + OMPTGT #pragma omp F S collapse(3) for (i = 0; i < 10; i++) for (int j = 10; j < 8; j++) @@ -137,6 +157,7 @@ __attribute__((noinline, noclone)) void N(f11) (int n) { int i; + OMPTGT #pragma omp F S for (i = 20; i < n; i++) { @@ -150,6 +171,7 @@ __attribute__((noinline, noclone)) void N(f12) (int n) { SC int i; + OMPTGT #pragma omp F S collapse(3) for (i = 0; i < 10; i++) for (int j = n; j < 8; j++) @@ -165,6 +187,7 @@ __attribute__((noinline, noclone)) void N(f13) (void) { int *i; + OMPTGT #pragma omp F S for (i = a; i < &a[1500]; i++) i[0] += 2; @@ -174,6 +197,7 @@ __attribute__((noinline, noclone)) void N(f14) (void) { SC float *i; + OMPTGT #pragma omp F S collapse(3) for (i = &b[0][0][0]; i < &b[0][0][10]; i++) for (float *j = &b[0][15][0]; j > &b[0][0][0]; j -= 10) @@ -188,27 +212,34 @@ N(test) (void) int i, j, k; for (i = 0; i < 1500; i++) a[i] = i - 25; + OMPTO (a); N(f0) (); + OMPFROM (a); for (i = 0; i < 1500; i++) if (a[i] != i - 23) return 1; N(f1) (); + OMPFROM (a); for (i = 0; i < 1500; i++) if (a[i] != i - 25) return 1; N(f2) (); + OMPFROM (a); for (i = 0; i < 1500; i++) if (a[i] != i - 29) return 1; N(f3) (1500LL - 1 - 23 - 48, -1LL + 25 - 48, 1LL); + OMPFROM (a); for (i = 0; i < 1500; i++) if (a[i] != i - 22) return 1; N(f3) (1500LL - 1 - 23 - 48, 1500LL - 1, 7LL); + OMPFROM (a); for (i = 0; i < 1500; i++) if (a[i] != i - 22) return 1; N(f4) (); + OMPFROM (a); for (i = 0; i < 1500; i++) if (a[i] != i - 22) return 1; @@ -216,31 +247,37 @@ N(test) (void) for (j = 0; j < 15; j++) for (k = 0; k < 10; k++) b[i][j][k] = i - 2.5 + 1.5 * j - 1.5 * k; + OMPTO (b); N(f5) (0, 10, 0, 15, 0, 10, 1, 1, 1); + OMPFROM (b); for (i = 0; i < 10; i++) for (j = 0; j < 15; j++) for (k = 0; k < 10; k++) if (b[i][j][k] != i + 1.5 * j - 1.5 * k) return 1; N(f5) (0, 10, 30, 15, 0, 10, 4, 5, 6); + OMPFROM (b); for (i = 0; i < 10; i++) for (j = 0; j < 15; j++) for (k = 0; k < 10; k++) if (b[i][j][k] != i + 1.5 * j - 1.5 * k) return 1; N(f6) (9, -1, 29, 0, 9, -1, -1, -2, -1); + OMPFROM (b); for (i = 0; i < 10; i++) for (j = 0; j < 15; j++) for (k = 0; k < 10; k++) if (b[i][j][k] != i - 4.5 + 1.5 * j - 1.5 * k) return 1; N(f7) (); + OMPFROM (b); for (i = 0; i < 10; i++) for (j = 0; j < 15; j++) for (k = 0; k < 10; k++) if (b[i][j][k] != i + 1.0 + 1.5 * j - 1.5 * k) return 1; N(f8) (); + OMPFROM (b); for (i = 0; i < 10; i++) for (j = 0; j < 15; j++) for (k = 0; k < 10; k++) @@ -250,6 +287,8 @@ N(test) (void) N(f10) (); N(f11) (10); N(f12) (12); + OMPFROM (a); + OMPFROM (b); for (i = 0; i < 1500; i++) if (a[i] != i - 22) return 1; @@ -260,6 +299,8 @@ N(test) (void) return 1; N(f13) (); N(f14) (); + OMPFROM (a); + OMPFROM (b); for (i = 0; i < 1500; i++) if (a[i] != i - 20) return 1; diff --git a/libgomp/testsuite/libgomp.c/for-4.c b/libgomp/testsuite/libgomp.c/for-4.c new file mode 100644 index 0000000..ef5465e --- /dev/null +++ b/libgomp/testsuite/libgomp.c/for-4.c @@ -0,0 +1,42 @@ +/* { dg-options "-std=gnu99 -fopenmp" } */ + +extern void abort (void); + +#define M(x, y, z) O(x, y, z) +#define O(x, y, z) x ## _ ## y ## _ ## z + +#define F taskloop +#define G taskloop +#define S +#define N(x) M(x, G, normal) +#include "for-2.h" +#undef S +#undef N +#undef F +#undef G + +#define F taskloop simd +#define G taskloop_simd +#define S +#define N(x) M(x, G, normal) +#include "for-2.h" +#undef S +#undef N +#undef F +#undef G + +int +main () +{ + int err = 0; + #pragma omp parallel reduction(|:err) + #pragma omp single + { + if (test_taskloop_normal () + || test_taskloop_simd_normal ()) + err = 1; + } + if (err) + abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/for-5.c b/libgomp/testsuite/libgomp.c/for-5.c new file mode 100644 index 0000000..84e636a --- /dev/null +++ b/libgomp/testsuite/libgomp.c/for-5.c @@ -0,0 +1,154 @@ +/* { dg-additional-options "-std=gnu99" } */ + +extern void abort (); + +#define M(x, y, z) O(x, y, z) +#define O(x, y, z) x ## _ ## y ## _ ## z + +#pragma omp declare target + +#define F for +#define G f +#define S +#define N(x) M(x, G, normal) +#include "for-2.h" +#undef S +#undef N +#undef F +#undef G + +#pragma omp end declare target + +#undef OMPFROM +#undef OMPTO +#define DO_PRAGMA(x) _Pragma (#x) +#define OMPFROM(v) DO_PRAGMA (omp target update from(v)) +#define OMPTO(v) DO_PRAGMA (omp target update to(v)) + +#define F target parallel for +#define G tpf +#include "for-1.h" +#undef F +#undef G + +#define F target simd +#define G t_simd +#define S +#define N(x) M(x, G, normal) +#include "for-2.h" +#undef S +#undef N +#undef F +#undef G + +#define F target parallel for simd +#define G tpf_simd +#include "for-1.h" +#undef F +#undef G + +#define F target teams distribute +#define G ttd +#define S +#define N(x) M(x, G, normal) +#include "for-2.h" +#undef S +#undef N +#undef F +#undef G + +#define F target teams distribute +#define G ttd_ds128 +#define S dist_schedule(static, 128) +#define N(x) M(x, G, normal) +#include "for-2.h" +#undef S +#undef N +#undef F +#undef G + +#define F target teams distribute simd +#define G ttds +#define S +#define N(x) M(x, G, normal) +#include "for-2.h" +#undef S +#undef N +#undef F +#undef G + +#define F target teams distribute simd +#define G ttds_ds128 +#define S dist_schedule(static, 128) +#define N(x) M(x, G, normal) +#include "for-2.h" +#undef S +#undef N +#undef F +#undef G + +#define F target teams distribute parallel for +#define G ttdpf +#include "for-1.h" +#undef F +#undef G + +#define F target teams distribute parallel for dist_schedule(static, 128) +#define G ttdpf_ds128 +#include "for-1.h" +#undef F +#undef G + +#define F target teams distribute parallel for simd +#define G ttdpfs +#include "for-1.h" +#undef F +#undef G + +#define F target teams distribute parallel for simd dist_schedule(static, 128) +#define G ttdpfs_ds128 +#include "for-1.h" +#undef F +#undef G + +int +main () +{ + if (test_tpf_static () + || test_tpf_static32 () + || test_tpf_auto () + || test_tpf_guided32 () + || test_tpf_runtime () + || test_t_simd_normal () + || test_tpf_simd_static () + || test_tpf_simd_static32 () + || test_tpf_simd_auto () + || test_tpf_simd_guided32 () + || test_tpf_simd_runtime () + || test_ttd_normal () + || test_ttd_ds128_normal () + || test_ttds_normal () + || test_ttds_ds128_normal () + || test_ttdpf_static () + || test_ttdpf_static32 () + || test_ttdpf_auto () + || test_ttdpf_guided32 () + || test_ttdpf_runtime () + || test_ttdpf_ds128_static () + || test_ttdpf_ds128_static32 () + || test_ttdpf_ds128_auto () + || test_ttdpf_ds128_guided32 () + || test_ttdpf_ds128_runtime () + || test_ttdpfs_static () + || test_ttdpfs_static32 () + || test_ttdpfs_auto () + || test_ttdpfs_guided32 () + || test_ttdpfs_runtime () + || test_ttdpfs_ds128_static () + || test_ttdpfs_ds128_static32 () + || test_ttdpfs_ds128_auto () + || test_ttdpfs_ds128_guided32 () + || test_ttdpfs_ds128_runtime ()) + abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/for-6.c b/libgomp/testsuite/libgomp.c/for-6.c new file mode 100644 index 0000000..7f3c65e --- /dev/null +++ b/libgomp/testsuite/libgomp.c/for-6.c @@ -0,0 +1,123 @@ +/* { dg-additional-options "-std=gnu99" } */ + +extern void abort (); + +#define M(x, y, z) O(x, y, z) +#define O(x, y, z) x ## _ ## y ## _ ## z + +#pragma omp declare target + +#define F for +#define G f +#define S +#define N(x) M(x, G, normal) +#include "for-2.h" +#undef S +#undef N +#undef F +#undef G + +#pragma omp end declare target + +#undef OMPTGT +#undef OMPFROM +#undef OMPTO +#define DO_PRAGMA(x) _Pragma (#x) +#define OMPTGT DO_PRAGMA (omp target) +#define OMPFROM(v) DO_PRAGMA (omp target update from(v)) +#define OMPTO(v) DO_PRAGMA (omp target update to(v)) + +#define F teams distribute +#define G td +#define S +#define N(x) M(x, G, normal) +#include "for-2.h" +#undef S +#undef N +#undef F +#undef G + +#define F teams distribute +#define G td_ds128 +#define S dist_schedule(static, 128) +#define N(x) M(x, G, normal) +#include "for-2.h" +#undef S +#undef N +#undef F +#undef G + +#define F teams distribute simd +#define G tds +#define S +#define N(x) M(x, G, normal) +#include "for-2.h" +#undef S +#undef N +#undef F +#undef G + +#define F teams distribute simd +#define G tds_ds128 +#define S dist_schedule(static, 128) +#define N(x) M(x, G, normal) +#include "for-2.h" +#undef S +#undef N +#undef F +#undef G + +#define F teams distribute parallel for +#define G tdpf +#include "for-1.h" +#undef F +#undef G + +#define F teams distribute parallel for dist_schedule(static, 128) +#define G tdpf_ds128 +#include "for-1.h" +#undef F +#undef G + +#define F teams distribute parallel for simd +#define G tdpfs +#include "for-1.h" +#undef F +#undef G + +#define F teams distribute parallel for simd dist_schedule(static, 128) +#define G tdpfs_ds128 +#include "for-1.h" +#undef F +#undef G + +int +main () +{ + if (test_td_normal () + || test_td_ds128_normal () + || test_tds_normal () + || test_tds_ds128_normal () + || test_tdpf_static () + || test_tdpf_static32 () + || test_tdpf_auto () + || test_tdpf_guided32 () + || test_tdpf_runtime () + || test_tdpf_ds128_static () + || test_tdpf_ds128_static32 () + || test_tdpf_ds128_auto () + || test_tdpf_ds128_guided32 () + || test_tdpf_ds128_runtime () + || test_tdpfs_static () + || test_tdpfs_static32 () + || test_tdpfs_auto () + || test_tdpfs_guided32 () + || test_tdpfs_runtime () + || test_tdpfs_ds128_static () + || test_tdpfs_ds128_static32 () + || test_tdpfs_ds128_auto () + || test_tdpfs_ds128_guided32 () + || test_tdpfs_ds128_runtime ()) + abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/linear-1.c b/libgomp/testsuite/libgomp.c/linear-1.c new file mode 100644 index 0000000..f86fb33 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/linear-1.c @@ -0,0 +1,250 @@ +int a[256]; + +__attribute__((noinline, noclone)) int +f1 (int i) +{ + #pragma omp parallel for linear (i: 4) + for (int j = 16; j < 64; j++) + { + a[i] = j; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) short int +f2 (short int i, char k) +{ + #pragma omp parallel for linear (i: k + 1) + for (long j = 16; j < 64; j++) + { + a[i] = j; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) long long int +f3 (long long int i, long long int k) +{ + #pragma omp parallel for linear (i: k) + for (short j = 16; j < 64; j++) + { + a[i] = j; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) int +f4 (int i) +{ + #pragma omp parallel for linear (i: 4) schedule(static, 3) + for (int j = 16; j < 64; j++) + { + a[i] = j; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) short int +f5 (short int i, char k) +{ + #pragma omp parallel for linear (i: k + 1) schedule(static, 5) + for (long j = 16; j < 64; j++) + { + a[i] = j; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) long long int +f6 (long long int i, long long int k) +{ + #pragma omp parallel for linear (i: k) schedule(static, 7) + for (short j = 16; j < 64; j++) + { + a[i] = j; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) int +f7 (int i) +{ + #pragma omp parallel for linear (i: 4) schedule(dynamic, 3) + for (int j = 16; j < 64; j++) + { + a[i] = j; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) short int +f8 (short int i, char k) +{ + #pragma omp parallel for linear (i: k + 1) schedule(dynamic, 5) + for (long j = 16; j < 64; j++) + { + a[i] = j; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) long long int +f9 (long long int i, long long int k) +{ + #pragma omp parallel for linear (i: k) schedule(dynamic, 7) + for (short j = 16; j < 64; j++) + { + a[i] = j; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) int +f10 (int i, long step) +{ + #pragma omp parallel for linear (i: 4) + for (int j = 16; j < 112; j += step) + { + a[i] = j / 2 + 8; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) short int +f11 (short int i, char k, char step) +{ + #pragma omp parallel for linear (i: k + 1) + for (long j = 16; j < 112; j += step) + { + a[i] = j / 2 + 8; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) long long int +f12 (long long int i, long long int k, int step) +{ + #pragma omp parallel for linear (i: k) + for (short j = 16; j < 112; j += step) + { + a[i] = j / 2 + 8; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) int +f13 (int i, long long int step) +{ + #pragma omp parallel for linear (i: 4) schedule(static, 3) + for (int j = 16; j < 112; j += step) + { + a[i] = j / 2 + 8; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) short int +f14 (short int i, char k, int step) +{ + #pragma omp parallel for linear (i: k + 1) schedule(static, 5) + for (long j = 16; j < 112; j += step) + { + a[i] = j / 2 + 8; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) long long int +f15 (long long int i, long long int k, long int step) +{ + #pragma omp parallel for linear (i: k) schedule(static, 7) + for (short j = 16; j < 112; j += step) + { + a[i] = j / 2 + 8; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) int +f16 (int i, long long int step) +{ + #pragma omp parallel for linear (i: 4) schedule(dynamic, 3) + for (int j = 16; j < 112; j += step) + { + a[i] = j / 2 + 8; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) short int +f17 (short int i, char k, int step) +{ + #pragma omp parallel for linear (i: k + 1) schedule(dynamic, 5) + for (long j = 16; j < 112; j += step) + { + a[i] = j / 2 + 8; + i += 4; + } + return i; +} + +__attribute__((noinline, noclone)) long long int +f18 (long long int i, long long int k, long int step) +{ + #pragma omp parallel for linear (i: k) schedule(dynamic, 7) + for (short j = 16; j < 112; j += step) + { + a[i] = j / 2 + 8; + i += 4; + } + return i; +} + +int +main () +{ +#define TEST(x) \ + if (x != 8 + 48 * 4) \ + __builtin_abort (); \ + for (int i = 0; i < 256; i++) \ + if (a[i] != (((i & 3) == 0 && i >= 8 \ + && i < 8 + 48 * 4) \ + ? ((i - 8) / 4) + 16 : 0)) \ + __builtin_abort (); \ + __builtin_memset (a, 0, sizeof (a)) + TEST (f1 (8)); + TEST (f2 (8, 3)); + TEST (f3 (8LL, 4LL)); + TEST (f4 (8)); + TEST (f5 (8, 3)); + TEST (f6 (8LL, 4LL)); + TEST (f7 (8)); + TEST (f8 (8, 3)); + TEST (f9 (8LL, 4LL)); + TEST (f10 (8, 2)); + TEST (f11 (8, 3, 2)); + TEST (f12 (8LL, 4LL, 2)); + TEST (f13 (8, 2)); + TEST (f14 (8, 3, 2)); + TEST (f15 (8LL, 4LL, 2)); + TEST (f16 (8, 2)); + TEST (f17 (8, 3, 2)); + TEST (f18 (8LL, 4LL, 2)); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/ordered-4.c b/libgomp/testsuite/libgomp.c/ordered-4.c new file mode 100644 index 0000000..8412d47 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/ordered-4.c @@ -0,0 +1,83 @@ +extern +#ifdef __cplusplus +"C" +#endif +void abort (void); + +void +foo (int i, char *j) +{ + #pragma omp atomic + j[i]++; + #pragma omp ordered threads + { + int t; + #pragma omp atomic read + t = j[i]; + if (t != 3) + abort (); + if (i > 1) + { + #pragma omp atomic read + t = j[i - 1]; + if (t == 2) + abort (); + } + if (i < 127) + { + #pragma omp atomic read + t = j[i + 1]; + if (t == 4) + abort (); + } + } + #pragma omp atomic + j[i]++; +} + +int +main () +{ + int i; + char j[128]; + #pragma omp parallel + { + #pragma omp for + for (i = 0; i < 128; i++) + j[i] = 0; + #pragma omp for ordered schedule(dynamic, 1) + for (i = 0; i < 128; i++) + { + #pragma omp atomic + j[i]++; + #pragma omp ordered threads + { + int t; + #pragma omp atomic read + t = j[i]; + if (t != 1) + abort (); + if (i > 1) + { + #pragma omp atomic read + t = j[i - 1]; + if (t == 0) + abort (); + } + if (i < 127) + { + #pragma omp atomic read + t = j[i + 1]; + if (t == 2) + abort (); + } + } + #pragma omp atomic + j[i]++; + } + #pragma omp for ordered schedule(static, 1) + for (i = 0; i < 128; i++) + foo (i, j); + } + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/pr66199-2.c b/libgomp/testsuite/libgomp.c/pr66199-2.c index ddb79de..2fc9eec 100644 --- a/libgomp/testsuite/libgomp.c/pr66199-2.c +++ b/libgomp/testsuite/libgomp.c/pr66199-2.c @@ -18,12 +18,11 @@ __attribute__((noinline, noclone)) void f2 (long a, long b, long c) { long d, e; - #pragma omp target teams distribute parallel for simd default(none) firstprivate (a, b) shared(u, v, w) linear(d) linear(c:5) lastprivate(e) + #pragma omp target teams distribute parallel for simd default(none) firstprivate (a, b, c) shared(u, v, w) linear(d) lastprivate(e) for (d = a; d < b; d++) { u[d] = v[d] + w[d]; - c += 5; - e = c; + e = c + d * 5; } } diff --git a/libgomp/testsuite/libgomp.c/pr66199-3.c b/libgomp/testsuite/libgomp.c/pr66199-3.c new file mode 100644 index 0000000..fe0ccb4 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr66199-3.c @@ -0,0 +1,50 @@ +/* PR middle-end/66199 */ +/* { dg-do run } */ +/* { dg-options "-O2 -fopenmp" } */ + +int u[1024], v[1024], w[1024]; + +__attribute__((noinline, noclone)) long +f1 (long a, long b) +{ + long d; + #pragma omp parallel for lastprivate (d) default(none) firstprivate (a, b) shared(u, v, w) + for (d = a; d < b; d++) + u[d] = v[d] + w[d]; + return d; +} + +__attribute__((noinline, noclone)) long +f2 (long a, long b, long c) +{ + long d, e; + #pragma omp parallel for lastprivate (d) default(none) firstprivate (a, b) shared(u, v, w) linear(c:5) lastprivate(e) + for (d = a; d < b; d++) + { + u[d] = v[d] + w[d]; + c += 5; + e = c; + } + return d + c + e; +} + +__attribute__((noinline, noclone)) long +f3 (long a1, long b1, long a2, long b2) +{ + long d1, d2; + #pragma omp parallel for default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) collapse(2) + for (d1 = a1; d1 < b1; d1++) + for (d2 = a2; d2 < b2; d2++) + u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2]; + return d1 + d2; +} + +int +main () +{ + if (f1 (0, 1024) != 1024 + || f2 (0, 1024, 17) != 1024 + 2 * (17 + 5 * 1024) + || f3 (0, 32, 0, 32) != 64) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/pr66199-4.c b/libgomp/testsuite/libgomp.c/pr66199-4.c new file mode 100644 index 0000000..a9b1bb8 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr66199-4.c @@ -0,0 +1,58 @@ +/* PR middle-end/66199 */ +/* { dg-do run } */ +/* { dg-options "-O2 -fopenmp" } */ + +#pragma omp declare target +int u[1024], v[1024], w[1024]; +#pragma omp end declare target + +__attribute__((noinline, noclone)) void +f1 (long a, long b) +{ + long d; + #pragma omp target teams distribute parallel for default(none) firstprivate (a, b) shared(u, v, w) + for (d = a; d < b; d++) + u[d] = v[d] + w[d]; +} + +__attribute__((noinline, noclone)) void +f2 (long a, long b, long c) +{ + long d, e; + #pragma omp target teams distribute parallel for default(none) firstprivate (a, b, c) shared(u, v, w) lastprivate(d, e) + for (d = a; d < b; d++) + { + u[d] = v[d] + w[d]; + e = c + d * 5; + } +} + +__attribute__((noinline, noclone)) void +f3 (long a1, long b1, long a2, long b2) +{ + long d1, d2; + #pragma omp target teams distribute parallel for default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) collapse(2) + for (d1 = a1; d1 < b1; d1++) + for (d2 = a2; d2 < b2; d2++) + u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2]; +} + +__attribute__((noinline, noclone)) void +f4 (long a1, long b1, long a2, long b2) +{ + long d1, d2; + #pragma omp target teams distribute parallel for default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) collapse(2) + for (d1 = a1; d1 < b1; d1++) + for (d2 = a2; d2 < b2; d2++) + u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2]; +} + +int +main () +{ + f1 (0, 1024); + f2 (0, 1024, 17); + f3 (0, 32, 0, 32); + f4 (0, 32, 0, 32); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/reduction-10.c b/libgomp/testsuite/libgomp.c/reduction-10.c new file mode 100644 index 0000000..3c95ebd --- /dev/null +++ b/libgomp/testsuite/libgomp.c/reduction-10.c @@ -0,0 +1,105 @@ +struct A { int t; }; +struct B { char t; }; +struct C { unsigned long long t; }; +struct D { long t; }; +void +add (struct B *x, struct B *y) +{ + x->t += y->t; +} +void +zero (struct B *x) +{ + x->t = 0; +} +void +orit (struct C *x, struct C *y) +{ + y->t |= x->t; +} +#pragma omp declare reduction(+:struct A:omp_out.t += omp_in.t) +#pragma omp declare reduction(+:struct B:add (&omp_out, &omp_in)) initializer(zero (&omp_priv)) +#pragma omp declare reduction(*:struct A:omp_out.t *= omp_in.t) initializer(omp_priv = { 1 }) +#pragma omp declare reduction(|:struct C:orit (&omp_in, &omp_out)) +#pragma omp declare reduction(&:struct D:omp_out.t = omp_out.t & omp_in.t) initializer(omp_priv = { ~0L }) +#pragma omp declare reduction(maxb:short:omp_out = omp_in > omp_out ? omp_in : omp_out) initializer(omp_priv = -6) + +struct B z[10]; + +__attribute__((noinline, noclone)) void +foo (struct A (*x)[3][2], struct A *y, struct D w[1][2], int p1, long p2, long p3, int p4, + int p5, long p6, short p7) +{ + struct C a[p7 + 4]; + short b[p7]; + int i; + for (i = 0; i < p7 + 4; i++) + { + if (i < p7) + b[i] = -6; + a[i].t = 0; + } + #pragma omp parallel for reduction(+:x[0:p1 + 1][:p2], z[:p3]) \ + reduction(*:y[:p4]) reduction(|:a[:p5]) \ + reduction(&:w[0:p6 - 1][:p6]) reduction(maxb:b) + for (i = 0; i < 128; i++) + { + x[i / 64][i % 3][(i / 4) & 1].t += i; + if ((i & 15) == 1) + y[0].t *= 3; + if ((i & 31) == 2) + y[1].t *= 7; + if ((i & 63) == 3) + y[2].t *= 17; + z[i / 32].t += (i & 3); + if (i < 4) + z[i].t += i; + a[i / 32].t |= 1ULL << (i & 30); + w[0][i & 1].t &= ~(1L << (i / 17 * 3)); + if ((i % 79) > b[0]) + b[0] = i % 79; + if ((i % 13) > b[1]) + b[1] = i % 13; + if ((i % 23) > b[2]) + b[2] = i % 23; + if ((i % 85) > b[3]) + b[3] = i % 85; + if ((i % 192) > b[4]) + b[4] = i % 192; + } + for (i = 0; i < 9; i++) + if (a[i].t != (i < 4 ? 0x55555555ULL : 0)) + __builtin_abort (); + if (b[0] != 78 || b[1] != 12 || b[2] != 22 || b[3] != 84 || b[4] != 127) + __builtin_abort (); +} + +int +main () +{ + struct A a[4][3][2] = {}; + static int a2[4][3][2] = {{{ 0, 0 }, { 0, 0 }, { 0, 0 }}, + {{ 312, 381 }, { 295, 356 }, { 337, 335 }}, + {{ 1041, 975 }, { 1016, 1085 }, { 935, 1060 }}, + {{ 0, 0 }, { 0, 0 }, { 0, 0 }}}; + struct A y[5] = { { 0 }, { 1 }, { 1 }, { 1 }, { 0 } }; + int y2[5] = { 0, 6561, 2401, 289, 0 }; + char z2[10] = { 48, 49, 50, 51, 0, 0, 0, 0, 0, 0 }; + struct D w[1][2] = { { { ~0L }, { ~0L } } }; + foo (&a[1], y + 1, w, 1, 3L, 4L, 3, 4, 2L, 5); + int i, j, k; + for (i = 0; i < 4; i++) + for (j = 0; j < 3; j++) + for (k = 0; k < 2; k++) + if (a[i][j][k].t != a2[i][j][k]) + __builtin_abort (); + for (i = 0; i < 5; i++) + if (y[i].t != y2[i]) + __builtin_abort (); + for (i = 0; i < 10; i++) + if (z[i].t != z2[i]) + __builtin_abort (); + if (w[0][0].t != ~0x249249L || w[0][1].t != ~0x249249L) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/reduction-7.c b/libgomp/testsuite/libgomp.c/reduction-7.c new file mode 100644 index 0000000..347c26f --- /dev/null +++ b/libgomp/testsuite/libgomp.c/reduction-7.c @@ -0,0 +1,64 @@ +char z[10] = { 0 }; + +__attribute__((noinline, noclone)) void +foo (int (*x)[3][2], int *y, long w[1][2]) +{ + unsigned long long a[9] = {}; + short b[5] = {}; + int i; + #pragma omp parallel for reduction(+:x[0:2][:][0:2], z[:4]) \ + reduction(*:y[:3]) reduction(|:a[:4]) \ + reduction(&:w[0:1][:2]) reduction(max:b) + for (i = 0; i < 128; i++) + { + x[i / 64][i % 3][(i / 4) & 1] += i; + if ((i & 15) == 1) + y[0] *= 3; + if ((i & 31) == 2) + y[1] *= 7; + if ((i & 63) == 3) + y[2] *= 17; + z[i / 32] += (i & 3); + if (i < 4) + z[i] += i; + a[i / 32] |= 1ULL << (i & 30); + w[0][i & 1] &= ~(1L << (i / 17 * 3)); + if ((i % 79) > b[0]) + b[0] = i % 79; + if ((i % 13) > b[1]) + b[1] = i % 13; + if ((i % 23) > b[2]) + b[2] = i % 23; + if ((i % 85) > b[3]) + b[3] = i % 85; + if ((i % 192) > b[4]) + b[4] = i % 192; + } + for (i = 0; i < 9; i++) + if (a[i] != (i < 4 ? 0x55555555ULL : 0)) + __builtin_abort (); + if (b[0] != 78 || b[1] != 12 || b[2] != 22 || b[3] != 84 || b[4] != 127) + __builtin_abort (); +} + +int +main () +{ + int a[4][3][2] = {}; + static int a2[4][3][2] = {{{ 0, 0 }, { 0, 0 }, { 0, 0 }}, + {{ 312, 381 }, { 295, 356 }, { 337, 335 }}, + {{ 1041, 975 }, { 1016, 1085 }, { 935, 1060 }}, + {{ 0, 0 }, { 0, 0 }, { 0, 0 }}}; + int y[5] = { 0, 1, 1, 1, 0 }; + int y2[5] = { 0, 6561, 2401, 289, 0 }; + char z2[10] = { 48, 49, 50, 51, 0, 0, 0, 0, 0, 0 }; + long w[1][2] = { ~0L, ~0L }; + foo (&a[1], y + 1, w); + if (__builtin_memcmp (a, a2, sizeof (a)) + || __builtin_memcmp (y, y2, sizeof (y)) + || __builtin_memcmp (z, z2, sizeof (z)) + || w[0][0] != ~0x249249L + || w[0][1] != ~0x249249L) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/reduction-8.c b/libgomp/testsuite/libgomp.c/reduction-8.c new file mode 100644 index 0000000..f4ec03a --- /dev/null +++ b/libgomp/testsuite/libgomp.c/reduction-8.c @@ -0,0 +1,98 @@ +struct A { int t; }; +struct B { char t; }; +struct C { unsigned long long t; }; +struct D { long t; }; +void +add (struct B *x, struct B *y) +{ + x->t += y->t; +} +void +zero (struct B *x) +{ + x->t = 0; +} +void +orit (struct C *x, struct C *y) +{ + y->t |= x->t; +} +#pragma omp declare reduction(+:struct A:omp_out.t += omp_in.t) +#pragma omp declare reduction(+:struct B:add (&omp_out, &omp_in)) initializer(zero (&omp_priv)) +#pragma omp declare reduction(*:struct A:omp_out.t *= omp_in.t) initializer(omp_priv = { 1 }) +#pragma omp declare reduction(|:struct C:orit (&omp_in, &omp_out)) +#pragma omp declare reduction(&:struct D:omp_out.t = omp_out.t & omp_in.t) initializer(omp_priv = { ~0L }) +#pragma omp declare reduction(maxb:short:omp_out = omp_in > omp_out ? omp_in : omp_out) initializer(omp_priv = -6) + +struct B z[10]; + +__attribute__((noinline, noclone)) void +foo (struct A (*x)[3][2], struct A *y, struct D w[1][2]) +{ + struct C a[9] = {}; + short b[5] = {}; + int i; + #pragma omp parallel for reduction(+:x[0:2][:][0:2], z[:4]) \ + reduction(*:y[:3]) reduction(|:a[:4]) \ + reduction(&:w[0:1][:2]) reduction(maxb:b) + for (i = 0; i < 128; i++) + { + x[i / 64][i % 3][(i / 4) & 1].t += i; + if ((i & 15) == 1) + y[0].t *= 3; + if ((i & 31) == 2) + y[1].t *= 7; + if ((i & 63) == 3) + y[2].t *= 17; + z[i / 32].t += (i & 3); + if (i < 4) + z[i].t += i; + a[i / 32].t |= 1ULL << (i & 30); + w[0][i & 1].t &= ~(1L << (i / 17 * 3)); + if ((i % 79) > b[0]) + b[0] = i % 79; + if ((i % 13) > b[1]) + b[1] = i % 13; + if ((i % 23) > b[2]) + b[2] = i % 23; + if ((i % 85) > b[3]) + b[3] = i % 85; + if ((i % 192) > b[4]) + b[4] = i % 192; + } + for (i = 0; i < 9; i++) + if (a[i].t != (i < 4 ? 0x55555555ULL : 0)) + __builtin_abort (); + if (b[0] != 78 || b[1] != 12 || b[2] != 22 || b[3] != 84 || b[4] != 127) + __builtin_abort (); +} + +int +main () +{ + struct A a[4][3][2] = {}; + static int a2[4][3][2] = {{{ 0, 0 }, { 0, 0 }, { 0, 0 }}, + {{ 312, 381 }, { 295, 356 }, { 337, 335 }}, + {{ 1041, 975 }, { 1016, 1085 }, { 935, 1060 }}, + {{ 0, 0 }, { 0, 0 }, { 0, 0 }}}; + struct A y[5] = { { 0 }, { 1 }, { 1 }, { 1 }, { 0 } }; + int y2[5] = { 0, 6561, 2401, 289, 0 }; + char z2[10] = { 48, 49, 50, 51, 0, 0, 0, 0, 0, 0 }; + struct D w[1][2] = { { { ~0L }, { ~0L } } }; + foo (&a[1], y + 1, w); + int i, j, k; + for (i = 0; i < 4; i++) + for (j = 0; j < 3; j++) + for (k = 0; k < 2; k++) + if (a[i][j][k].t != a2[i][j][k]) + __builtin_abort (); + for (i = 0; i < 5; i++) + if (y[i].t != y2[i]) + __builtin_abort (); + for (i = 0; i < 10; i++) + if (z[i].t != z2[i]) + __builtin_abort (); + if (w[0][0].t != ~0x249249L || w[0][1].t != ~0x249249L) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/reduction-9.c b/libgomp/testsuite/libgomp.c/reduction-9.c new file mode 100644 index 0000000..13605c1 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/reduction-9.c @@ -0,0 +1,71 @@ +char z[10] = { 0 }; + +__attribute__((noinline, noclone)) void +foo (int (*x)[3][2], int *y, long w[1][2], int p1, long p2, long p3, int p4, + int p5, long p6, short p7) +{ + unsigned long long a[p7 + 4]; + short b[p7]; + int i; + for (i = 0; i < p7 + 4; i++) + { + if (i < p7) + b[i] = -6; + a[i] = 0; + } + #pragma omp parallel for reduction(+:x[0:p1 + 1][:p2], z[:p3]) \ + reduction(*:y[:p4]) reduction(|:a[:p5]) \ + reduction(&:w[0:p6 - 1][:p6]) reduction(max:b) + for (i = 0; i < 128; i++) + { + x[i / 64][i % 3][(i / 4) & 1] += i; + if ((i & 15) == 1) + y[0] *= 3; + if ((i & 31) == 2) + y[1] *= 7; + if ((i & 63) == 3) + y[2] *= 17; + z[i / 32] += (i & 3); + if (i < 4) + z[i] += i; + a[i / 32] |= 1ULL << (i & 30); + w[0][i & 1] &= ~(1L << (i / 17 * 3)); + if ((i % 79) > b[0]) + b[0] = i % 79; + if ((i % 13) > b[1]) + b[1] = i % 13; + if ((i % 23) > b[2]) + b[2] = i % 23; + if ((i % 85) > b[3]) + b[3] = i % 85; + if ((i % 192) > b[4]) + b[4] = i % 192; + } + for (i = 0; i < 9; i++) + if (a[i] != (i < 4 ? 0x55555555ULL : 0)) + __builtin_abort (); + if (b[0] != 78 || b[1] != 12 || b[2] != 22 || b[3] != 84 || b[4] != 127) + __builtin_abort (); +} + +int +main () +{ + int a[4][3][2] = {}; + static int a2[4][3][2] = {{{ 0, 0 }, { 0, 0 }, { 0, 0 }}, + {{ 312, 381 }, { 295, 356 }, { 337, 335 }}, + {{ 1041, 975 }, { 1016, 1085 }, { 935, 1060 }}, + {{ 0, 0 }, { 0, 0 }, { 0, 0 }}}; + int y[5] = { 0, 1, 1, 1, 0 }; + int y2[5] = { 0, 6561, 2401, 289, 0 }; + char z2[10] = { 48, 49, 50, 51, 0, 0, 0, 0, 0, 0 }; + long w[1][2] = { ~0L, ~0L }; + foo (&a[1], y + 1, w, 1, 3L, 4L, 3, 4, 2L, 5); + if (__builtin_memcmp (a, a2, sizeof (a)) + || __builtin_memcmp (y, y2, sizeof (y)) + || __builtin_memcmp (z, z2, sizeof (z)) + || w[0][0] != ~0x249249L + || w[0][1] != ~0x249249L) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-1.c b/libgomp/testsuite/libgomp.c/target-1.c index f734d3c2..c7abb00 100644 --- a/libgomp/testsuite/libgomp.c/target-1.c +++ b/libgomp/testsuite/libgomp.c/target-1.c @@ -34,7 +34,7 @@ fn2 (int x, int y, int z) fn1 (b, c, x); #pragma omp target data map(to: b) { - #pragma omp target map(tofrom: c) + #pragma omp target map(tofrom: c, s) #pragma omp teams num_teams(y) thread_limit(z) reduction(+:s) firstprivate(x) #pragma omp distribute dist_schedule(static, 4) collapse(1) for (j=0; j < x; j += y) @@ -52,7 +52,7 @@ fn3 (int x) double b[1024], c[1024], s = 0; int i; fn1 (b, c, x); - #pragma omp target map(to: b, c) + #pragma omp target map(to: b, c) map(tofrom:s) #pragma omp parallel for reduction(+:s) for (i = 0; i < x; i++) tgt (), s += b[i] * c[i]; @@ -66,7 +66,8 @@ fn4 (int x, double *p) int i; fn1 (b, c, x); fn1 (d + x, p + x, x); - #pragma omp target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)]) + #pragma omp target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)]) \ + map(tofrom: s) #pragma omp parallel for reduction(+:s) for (i = 0; i < x; i++) s += b[i] * c[i] + d[x + i] + p[x + i]; diff --git a/libgomp/testsuite/libgomp.c/target-11.c b/libgomp/testsuite/libgomp.c/target-11.c new file mode 100644 index 0000000..625c286 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-11.c @@ -0,0 +1,86 @@ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +#include <stdlib.h> +#include <assert.h> + +#define N 32 + +void test_array_section (int *p) +{ + #pragma omp target data map(alloc: p[0:N]) + { + int ok = 1; + for (int i = 10; i < 10 + 4; i++) + p[i] = 997 * i; + + #pragma omp target map(always to:p[10:4]) map(tofrom: ok) + for (int i = 10; i < 10 + 4; i++) + if (p[i] != 997 * i) + ok = 0; + + assert (ok); + + #pragma omp target map(always from:p[7:9]) + for (int i = 0; i < N; i++) + p[i] = i; + } +} + +int main () +{ + int aa = 0, bb = 0, cc = 0, dd = 0; + + #pragma omp target data map(tofrom: aa) map(to: bb) map(from: cc, dd) + { + int ok; + aa = bb = cc = 1; + + /* Set dd on target to 0 for the further check. */ + #pragma omp target map(always to: dd) + ; + + dd = 1; + #pragma omp target map(tofrom: aa) map(always to: bb) \ + map(always from: cc) map(to: dd) map(from: ok) + { + /* bb is always to, aa and dd are not. */ + ok = (aa == 0) && (bb == 1) && (dd == 0); + aa = bb = cc = dd = 2; + } + + assert (ok); + assert (aa == 1); + assert (bb == 1); + assert (cc == 2); /* cc is always from. */ + assert (dd == 1); + + dd = 3; + #pragma omp target map(from: cc) map(always to: dd) map(from: ok) + { + ok = (dd == 3); /* dd is always to. */ + cc = dd = 4; + } + + assert (ok); + assert (cc == 2); + assert (dd == 3); + } + + assert (aa == 2); + assert (bb == 1); + assert (cc == 4); + assert (dd == 4); + + int *array = calloc (N, sizeof (int)); + test_array_section (array); + + for (int i = 0; i < 7; i++) + assert (array[i] == 0); + for (int i = 7; i < 7 + 9; i++) + assert (array[i] == i); + for (int i = 7 + 9; i < N; i++) + assert (array[i] == 0); + + free (array); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-12.c b/libgomp/testsuite/libgomp.c/target-12.c new file mode 100644 index 0000000..e6b0094 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-12.c @@ -0,0 +1,130 @@ +#include <omp.h> +#include <stdlib.h> + +int +main () +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + int err; + int q[128], i; + void *p; + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + for (i = 0; i < 128; i++) + q[i] = i; + + p = omp_target_alloc (130 * sizeof (int), d); + if (p == NULL) + return 0; + + if (omp_target_memcpy_rect (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL, NULL, + d, id) < 3 + || omp_target_memcpy_rect (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL, + NULL, id, d) < 3 + || omp_target_memcpy_rect (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL, + NULL, id, id) < 3) + abort (); + + if (omp_target_associate_ptr (q, p, 128 * sizeof (int), sizeof (int), d) == 0) + { + size_t volume[3] = { 128, 0, 0 }; + size_t dst_offsets[3] = { 0, 0, 0 }; + size_t src_offsets[3] = { 1, 0, 0 }; + size_t dst_dimensions[3] = { 128, 0, 0 }; + size_t src_dimensions[3] = { 128, 0, 0 }; + + if (omp_target_associate_ptr (q, p, 128 * sizeof (int), sizeof (int), d) != 0) + abort (); + + if (omp_target_is_present (q, d) != 1 + || omp_target_is_present (&q[32], d) != 1 + || omp_target_is_present (&q[128], d) != 1) + abort (); + + if (omp_target_memcpy (p, q, 128 * sizeof (int), sizeof (int), 0, + d, id) != 0) + abort (); + + #pragma omp target if (d >= 0) device (d >= 0 ? d : 0) map(alloc:q[0:32]) map(from:err) + { + int j; + err = 0; + for (j = 0; j < 128; j++) + if (q[j] != j) + err = 1; + else + q[j] += 4; + } + + if (err) + abort (); + + if (omp_target_memcpy_rect (q, p, sizeof (int), 1, volume, + dst_offsets, src_offsets, dst_dimensions, + src_dimensions, id, d) != 0) + abort (); + + for (i = 0; i < 128; i++) + if (q[i] != i + 4) + abort (); + + volume[2] = 2; + volume[1] = 3; + volume[0] = 6; + dst_offsets[2] = 1; + dst_offsets[1] = 0; + dst_offsets[0] = 0; + src_offsets[2] = 1; + src_offsets[1] = 0; + src_offsets[0] = 3; + dst_dimensions[2] = 2; + dst_dimensions[1] = 3; + dst_dimensions[0] = 6; + src_dimensions[2] = 3; + src_dimensions[1] = 4; + src_dimensions[0] = 6; + if (omp_target_memcpy_rect (p, q, sizeof (int), 3, volume, + dst_offsets, src_offsets, dst_dimensions, + src_dimensions, d, id) != 0) + abort (); + + #pragma omp target if (d >= 0) device (d >= 0 ? d : 0) map(alloc:q[0:32]) map(from:err) + { + int j, k, l; + err = 0; + for (j = 0; j < 6; j++) + for (k = 0; k < 3; k++) + for (l = 0; l < 2; l++) + if (q[j * 6 + k * 2 + l] != 3 * 12 + 4 + 1 + l + k * 3 + j * 12) + err = 1; + } + + if (err) + abort (); + + if (omp_target_memcpy (p, p, 10 * sizeof (int), 51 * sizeof (int), + 111 * sizeof (int), d, d) != 0) + abort (); + + #pragma omp target if (d >= 0) device (d >= 0 ? d : 0) map(alloc:q[0:32]) map(from:err) + { + int j; + err = 0; + for (j = 0; j < 10; j++) + if (q[50 + j] != q[110 + j]) + err = 1; + } + + if (err) + abort (); + + if (omp_target_disassociate_ptr (q, d) != 0) + abort (); + } + + omp_target_free (p, d); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-13.c b/libgomp/testsuite/libgomp.c/target-13.c new file mode 100644 index 0000000..168850b --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-13.c @@ -0,0 +1,45 @@ +#ifdef __cplusplus +extern "C" +#else +extern +#endif +void abort (void); +struct S { int s, t; }; + +void +foo () +{ + int x = 5, y = 6, err = 0; + struct S u = { 7, 8 }, v = { 9, 10 }; + double s = 11.5, t = 12.5; + #pragma omp target private (x, u, s) firstprivate (y, v, t) map(from:err) + { + x = y; + u = v; + s = t; + err = (x != 6 || y != 6 + || u.s != 9 || u.t != 10 || v.s != 9 || v.t != 10 + || s != 12.5 || t != 12.5); + x += 1; + y += 2; + u.s += 3; + v.t += 4; + s += 2.5; + t += 3.0; + if (x != 7 || y != 8 + || u.s != 12 || u.t != 10 || v.s != 9 || v.t != 14 + || s != 15.0 || t != 15.5) + err = 1; + } + if (err || x != 5 || y != 6 + || u.s != 7 || u.t != 8 || v.s != 9 || v.t != 10 + || s != 11.5 || t != 12.5) + abort (); +} + +int +main () +{ + foo (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-14.c b/libgomp/testsuite/libgomp.c/target-14.c new file mode 100644 index 0000000..17d3834 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-14.c @@ -0,0 +1,38 @@ +#include <omp.h> +#include <stdlib.h> + +int +main () +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + int err; + void *p; + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + p = omp_target_alloc (128 * sizeof (int), d); + if (p == NULL) + return 0; + + #pragma omp target is_device_ptr (p) if (d >= 0) device (d >= 0 ? d : 0) + { + int i, *q = (int *) p; + for (i = 0; i < 128; i++) + q[i] = i + 7; + } + #pragma omp target is_device_ptr (p) if (d >= 0) device (d >= 0 ? d : 0) map(from:err) + { + int i; + err = 0; + for (i = 0; i < 128; i++) + if (((int *) p)[i] != i + 7) + err = 1; + } + if (err) + abort (); + + omp_target_free (p, d); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-15.c b/libgomp/testsuite/libgomp.c/target-15.c new file mode 100644 index 0000000..fee9252 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-15.c @@ -0,0 +1,74 @@ +extern void abort (void); + +void +foo (int *x) +{ + int a[10], b[15], err, i; + for (i = 0; i < 10; i++) + a[i] = 7 * i; + for (i = 0; i < 15; i++) + b[i] = 8 * i; + #pragma omp target map(to:x[5:10], a[0:10], b[5:10]) map(from:err) + { + err = 0; + for (i = 0; i < 10; i++) + if (x[5 + i] != 20 + 4 * i + || a[i] != 7 * i + || b[5 + i] != 40 + 8 * i) + err = 1; + } + if (err) + abort (); +} + +void +bar (int n, int v) +{ + int a[n], b[n], c[n], d[n], e[n], err, i; + int (*x)[n] = &c; + for (i = 0; i < n; i++) + { + (*x)[i] = 4 * i; + a[i] = 7 * i; + b[i] = 8 * i; + } + #pragma omp target map(to:x[0][5:10], a[0:10], b[5:10]) map(from:err) + { + err = 0; + for (i = 0; i < 10; i++) + if ((*x)[5 + i] != 20 + 4 * i + || a[i] != 7 * i + || b[5 + i] != 40 + 8 * i) + err = 1; + } + if (err) + abort (); + for (i = 0; i < n; i++) + { + (*x)[i] = 9 * i; + a[i] = 12 * i; + b[i] = 13 * i; + } + #pragma omp target map(to:x[0][v:v+5], a[v-5:v+5], b[v:v+5]) map(from:err) + { + err = 0; + for (i = 0; i < 10; i++) + if ((*x)[5 + i] != 45 + 9 * i + || a[i] != 12 * i + || b[5 + i] != 65 + 13 * i) + err = 1; + } + if (err) + abort (); +} + +int +main () +{ + int x[15], i; + for (i = 0; i < 15; i++) + x[i] = 4 * i; + foo (x); + bar (15, 5); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-16.c b/libgomp/testsuite/libgomp.c/target-16.c new file mode 100644 index 0000000..7b0919b --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-16.c @@ -0,0 +1,45 @@ +extern void abort (void); + +void +foo (int n) +{ + int a[n], i, err; + for (i = 0; i < n; i++) + a[i] = 7 * i; + #pragma omp target firstprivate (a) map(from:err) private (i) + { + err = 0; + for (i = 0; i < n; i++) + if (a[i] != 7 * i) + err = 1; + } + if (err) + abort (); +} + +void +bar (int n) +{ + int a[n], i, err; + #pragma omp target private (a) map(from:err) + { + #pragma omp parallel for + for (i = 0; i < n; i++) + a[i] = 7 * i; + err = 0; + #pragma omp parallel for reduction(|:err) + for (i = 0; i < n; i++) + if (a[i] != 7 * i) + err |= 1; + } + if (err) + abort (); +} + +int +main () +{ + foo (7); + bar (7); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-17.c b/libgomp/testsuite/libgomp.c/target-17.c new file mode 100644 index 0000000..4a76201 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-17.c @@ -0,0 +1,99 @@ +extern void abort (void); + +void +foo (int n) +{ + int a[n], i, err; + for (i = 0; i < n; i++) + a[i] = 5 * i; + #pragma omp target map(to:a) map(from:err) private(i) + { + err = 0; + for (i = 0; i < n; i++) + if (a[i] != 5 * i) + err = 1; + } + if (err) + abort (); + for (i = 0; i < n; i++) + a[i] += i; + #pragma omp target map(from:err) private(i) + { + err = 0; + for (i = 0; i < n; i++) + if (a[i] != 6 * i) + err = 1; + } + if (err) + abort (); + for (i = 0; i < n; i++) + a[i] += i; + #pragma omp target firstprivate (a) map(from:err) private(i) + { + err = 0; + for (i = 0; i < n; i++) + if (a[i] != 7 * i) + err = 1; + } + if (err) + abort (); + int on = n; + #pragma omp target firstprivate (n) map(tofrom: n) + { + n++; + } + if (on != n) + abort (); + #pragma omp target map(tofrom: n) private (n) + { + n = 25; + } + if (on != n) + abort (); + for (i = 0; i < n; i++) + a[i] += i; + #pragma omp target map(to:a) firstprivate (a) map(from:err) private(i) + { + err = 0; + for (i = 0; i < n; i++) + if (a[i] != 8 * i) + err = 1; + } + if (err) + abort (); + for (i = 0; i < n; i++) + a[i] += i; + #pragma omp target firstprivate (a) map(to:a) map(from:err) private(i) + { + err = 0; + for (i = 0; i < n; i++) + if (a[i] != 9 * i) + err = 1; + } + if (err) + abort (); + for (i = 0; i < n; i++) + a[i] += i; + #pragma omp target map(tofrom:a) map(from:err) private(a, i) + { + err = 0; + for (i = 0; i < n; i++) + a[i] = 7; + #pragma omp parallel for reduction(|:err) + for (i = 0; i < n; i++) + if (a[i] != 7) + err |= 1; + } + if (err) + abort (); + for (i = 0; i < n; i++) + if (a[i] != 10 * i) + abort (); +} + +int +main () +{ + foo (9); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-18.c b/libgomp/testsuite/libgomp.c/target-18.c new file mode 100644 index 0000000..cbacaf6 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-18.c @@ -0,0 +1,52 @@ +extern void abort (void); + +void +foo (int n) +{ + int a[4] = { 0, 1, 2, 3 }, b[n]; + int *p = a + 1, i, err; + for (i = 0; i < n; i++) + b[i] = 9 + i; + #pragma omp target data map(to:a) + #pragma omp target data use_device_ptr(p) map(from:err) + #pragma omp target is_device_ptr(p) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 4; i++) + if (p[i - 1] != i) + err = 1; + } + if (err) + abort (); + for (i = 0; i < 4; i++) + a[i] = 23 + i; + #pragma omp target data map(to:a) + #pragma omp target data use_device_ptr(a) map(from:err) + #pragma omp target is_device_ptr(a) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 4; i++) + if (a[i] != 23 + i) + err = 1; + } + if (err) + abort (); + #pragma omp target data map(to:b) + #pragma omp target data use_device_ptr(b) map(from:err) + #pragma omp target is_device_ptr(b) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 4; i++) + if (b[i] != 9 + i) + err = 1; + } + if (err) + abort (); +} + +int +main () +{ + foo (9); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-19.c b/libgomp/testsuite/libgomp.c/target-19.c new file mode 100644 index 0000000..710c507 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-19.c @@ -0,0 +1,127 @@ +extern void abort (void); + +void +foo (int *p, int *q, int *r, int n, int m) +{ + int i, err, *s = r; + #pragma omp target data map(to:p[0:8]) + { + /* For zero length array sections, p points to the start of + already mapped range, q to the end of it, and r does not point + to an mapped range. */ + #pragma omp target map(alloc:p[:0]) map(to:q[:0]) map(from:r[:0]) private(i) map(from:err) firstprivate (s) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1 || q[i - 8] != i + 1) + err = 1; + if (p + 8 != q || (r != (int *) 0 && r != s)) + err = 1; + } + if (err) + abort (); + /* Implicit mapping of pointers behaves the same way. */ + #pragma omp target private(i) map(from:err) firstprivate (s) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1 || q[i - 8] != i + 1) + err = 1; + if (p + 8 != q || (r != (int *) 0 && r != s)) + err = 1; + } + if (err) + abort (); + /* And zero-length array sections, though not known at compile + time, behave the same. */ + #pragma omp target map(p[:n]) map(tofrom:q[:n]) map(alloc:r[:n]) private(i) map(from:err) firstprivate (s) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1 || q[i - 8] != i + 1) + err = 1; + if (p + 8 != q || (r != (int *) 0 && r != s)) + err = 1; + } + if (err) + abort (); + /* Non-zero length array sections, though not known at compile, + behave differently. */ + #pragma omp target map(p[:m]) map(tofrom:q[:m]) map(to:r[:m]) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (q[0] != 9 || r[0] != 10) + err = 1; + } + if (err) + abort (); + #pragma omp target data map(to:q[0:1]) + { + /* For zero length array sections, p points to the start of + already mapped range, q points to the start of another one, + and r to the end of the second one. */ + #pragma omp target map(to:p[:0]) map(from:q[:0]) map(tofrom:r[:0]) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (q[0] != 9 || r != q + 1) + err = 1; + } + if (err) + abort (); + /* Implicit mapping of pointers behaves the same way. */ + #pragma omp target private(i) map(from:err) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (q[0] != 9 || r != q + 1) + err = 1; + } + if (err) + abort (); + /* And zero-length array sections, though not known at compile + time, behave the same. */ + #pragma omp target map(p[:n]) map(alloc:q[:n]) map(from:r[:n]) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (q[0] != 9 || r != q + 1) + err = 1; + } + if (err) + abort (); + /* Non-zero length array sections, though not known at compile, + behave differently. */ + #pragma omp target map(p[:m]) map(alloc:q[:m]) map(tofrom:r[:m]) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (q[0] != 9 || r[0] != 10) + err = 1; + } + if (err) + abort (); + } + } +} + +int +main () +{ + int a[32], i; + for (i = 0; i < 32; i++) + a[i] = i; + foo (a + 1, a + 9, a + 10, 0, 1); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-2.c b/libgomp/testsuite/libgomp.c/target-2.c index ada8dad..0ba766c 100644 --- a/libgomp/testsuite/libgomp.c/target-2.c +++ b/libgomp/testsuite/libgomp.c/target-2.c @@ -23,7 +23,7 @@ fn2 (int x) int i; fn1 (b, c, x); fn1 (e, d + x, x); - #pragma omp target map(to: b, c[:x], d[x:x], e) + #pragma omp target map(to: b, c[:x], d[x:x], e) map(tofrom: s) #pragma omp parallel for reduction(+:s) for (i = 0; i < x; i++) s += b[i] * c[i] + d[x + i] + sizeof (b) - sizeof (c); @@ -38,7 +38,7 @@ fn3 (int x) int i; fn1 (b, c, x); fn1 (e, d, x); - #pragma omp target + #pragma omp target map(tofrom: s) #pragma omp parallel for reduction(+:s) for (i = 0; i < x; i++) s += b[i] * c[i] + d[i]; @@ -56,7 +56,7 @@ fn4 (int x) #pragma omp target data map(from: b, c[:x], d[x:x], e) { #pragma omp target update to(b, c[:x], d[x:x], e) - #pragma omp target map(c[:x], d[x:x]) + #pragma omp target map(c[:x], d[x:x], s) #pragma omp parallel for reduction(+:s) for (i = 0; i < x; i++) { diff --git a/libgomp/testsuite/libgomp.c/target-20.c b/libgomp/testsuite/libgomp.c/target-20.c new file mode 100644 index 0000000..3f4e798 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-20.c @@ -0,0 +1,120 @@ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +#include <stdlib.h> +#include <assert.h> + +#define N 40 + +int sum; +int var1 = 1; +int var2 = 2; + +#pragma omp declare target +int D[N]; +#pragma omp end declare target + +void enter_data (int *X) +{ + #pragma omp target enter data map(to: var1, var2, X[:N]) map(alloc: sum) +} + +void exit_data_0 (int *D) +{ + #pragma omp target exit data map(delete: D[:N]) +} + +void exit_data_1 () +{ + #pragma omp target exit data map(from: var1) +} + +void exit_data_2 (int *X) +{ + #pragma omp target exit data map(from: var2) map(release: X[:N], sum) +} + +void exit_data_3 (int *p) +{ + #pragma omp target exit data map(from: p[:0]) +} + +void test_nested () +{ + int X = 0, Y = 0, Z = 0; + + #pragma omp target data map(from: X, Y, Z) + { + #pragma omp target data map(from: X, Y, Z) + { + #pragma omp target map(from: X, Y, Z) + X = Y = Z = 1337; + assert (X == 0); + assert (Y == 0); + assert (Z == 0); + + #pragma omp target exit data map(from: X) map(release: Y) + assert (X == 0); + assert (Y == 0); + + #pragma omp target exit data map(release: Y) map(delete: Z) + assert (Y == 0); + assert (Z == 0); + } + assert (X == 1337); + assert (Y == 0); + assert (Z == 0); + + #pragma omp target map(from: X) + X = 2448; + assert (X == 2448); + assert (Y == 0); + assert (Z == 0); + + X = 4896; + } + assert (X == 4896); + assert (Y == 0); + assert (Z == 0); +} + +int main () +{ + int *X = malloc (N * sizeof (int)); + int *Y = malloc (N * sizeof (int)); + X[10] = 10; + Y[20] = 20; + enter_data (X); + + exit_data_0 (D); /* This should have no effect on D. */ + + #pragma omp target map(alloc: var1, var2, X[:N]) map(to: Y[:N]) map(always from: sum) + { + var1 += X[10]; + var2 += Y[20]; + sum = var1 + var2; + D[sum]++; + } + + assert (var1 == 1); + assert (var2 == 2); + assert (sum == 33); + + exit_data_1 (); + assert (var1 == 11); + assert (var2 == 2); + + /* Increase refcount of already mapped X[0:N]. */ + #pragma omp target enter data map(alloc: X[16:1]) + + exit_data_2 (X); + assert (var2 == 22); + + exit_data_3 (X + 5); /* Unmap X[0:N]. */ + + free (X); + free (Y); + + test_nested (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-21.c b/libgomp/testsuite/libgomp.c/target-21.c new file mode 100644 index 0000000..41498cf --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-21.c @@ -0,0 +1,79 @@ +extern void abort (void); +union U { int x; long long y; }; +struct T { int a; union U b; int c; }; +struct S { int s; int u; struct T v; int x[10]; union U w; int y[10]; int z[10]; }; +volatile int z; + +int +main () +{ + struct S s; + s.s = 0; + s.u = 1; + s.v.a = 2; + s.v.b.y = 3LL; + s.v.c = 19; + s.w.x = 4; + s.x[0] = 7; + s.x[1] = 8; + s.y[3] = 9; + s.y[4] = 10; + s.y[5] = 11; + int err = 0; + #pragma omp target map (to:s.v.b, s.u, s.x[0:z + 2]) \ + map (tofrom:s.y[3:3]) \ + map (from: s.w, s.z[z + 1:z + 3], err) + { + err = 0; + if (s.u != 1 || s.v.b.y != 3LL || s.x[0] != 7 || s.x[1] != 8 + || s.y[3] != 9 || s.y[4] != 10 || s.y[5] != 11) + err = 1; + s.w.x = 6; + s.y[3] = 12; + s.y[4] = 13; + s.y[5] = 14; + s.z[1] = 15; + s.z[2] = 16; + s.z[3] = 17; + } + if (err || s.w.x != 6 || s.y[3] != 12 || s.y[4] != 13 || s.y[5] != 14 + || s.z[1] != 15 || s.z[2] != 16 || s.z[3] != 17) + abort (); + s.u++; + s.v.a++; + s.v.b.y++; + s.w.x++; + s.x[1] = 18; + s.z[0] = 19; + #pragma omp target data map (tofrom: s) + #pragma omp target map (always to: s.w, s.x[1], err) map (alloc:s.u, s.v.b, s.z[z:z + 1]) + { + err = 0; + if (s.u != 2 || s.v.b.y != 4LL || s.w.x != 7 || s.x[1] != 18 || s.z[0] != 19) + err = 1; + s.w.x = 8; + s.x[1] = 20; + s.z[0] = 21; + } + if (err || s.w.x != 8 || s.x[1] != 20 || s.z[0] != 21) + abort (); + s.u++; + s.v.a++; + s.v.b.y++; + s.w.x++; + s.x[0] = 22; + s.x[1] = 23; + #pragma omp target data map (from: s.w, s.x[0:2]) map (to: s.v.b, s.u) + #pragma omp target map (always to: s.w, s.x[0:2], err) map (alloc:s.u, s.v.b) + { + err = 0; + if (s.u != 3 || s.v.b.y != 5LL || s.w.x != 9 || s.x[0] != 22 || s.x[1] != 23) + err = 1; + s.w.x = 11; + s.x[0] = 24; + s.x[1] = 25; + } + if (err || s.w.x != 11 || s.x[0] != 24 || s.x[1] != 25) + abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-22.c b/libgomp/testsuite/libgomp.c/target-22.c new file mode 100644 index 0000000..aad8a0a --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-22.c @@ -0,0 +1,51 @@ +extern void abort (void); +struct T { int a; int *b; int c; }; +struct S { int *s; char *u; struct T v; short *w; }; +volatile int z; + +int +main () +{ + struct S s; + int a[32], i; + char b[32]; + short c[32]; + for (i = 0; i < 32; i++) + { + a[i] = i; + b[i] = 32 + i; + c[i] = 64 + i; + } + s.s = a; + s.u = b + 2; + s.v.b = a + 16; + s.w = c + 3; + int err = 0; + #pragma omp target map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \ + map (tofrom:s.s[3:3]) \ + map (from: s.w[z:4], err) private (i) + { + err = 0; + for (i = 0; i < 7; i++) + if (s.v.b[i] != 16 + i) + err = 1; + for (i = 1; i < 5; i++) + if (s.u[i] != 34 + i) + err = 1; + for (i = 3; i < 6; i++) + if (s.s[i] != i) + err = 1; + else + s.s[i] = 128 + i; + for (i = 0; i < 4; i++) + s.w[i] = 96 + i; + } + if (err) + abort (); + for (i = 0; i < 32; i++) + if (a[i] != ((i >= 3 && i < 6) ? 128 + i : i) + || b[i] != 32 + i + || c[i] != ((i >= 3 && i < 7) ? 93 + i : 64 + i)) + abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-23.c b/libgomp/testsuite/libgomp.c/target-23.c new file mode 100644 index 0000000..fb1532a --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-23.c @@ -0,0 +1,48 @@ +extern void abort (void); +struct S { int s; int *u; int v[5]; }; +volatile int z; + +int +main () +{ + int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0; + struct S s = { 9, u + 3, { 10, 11, 12, 13, 14 } }; + int *v = u + 4; + #pragma omp target enter data map (to: s.s, s.u[0:5]) map (alloc: s.v[1:3]) + s.s++; + u[3]++; + s.v[1]++; + #pragma omp target update to (s.s) to (s.u[0:2], s.v[1:3]) + #pragma omp target map (alloc: s.s, s.v[1:3]) map (from: err) + { + err = 0; + if (s.s != 10 || s.v[1] != 12 || s.v[2] != 12 || s.v[3] != 13) + err = 1; + if (v[-1] != 4 || v[0] != 4 || v[1] != 5 || v[2] != 6 || v[3] != 7) + err = 1; + s.s++; + s.v[2] += 2; + v[-1] = 5; + v[3] = 9; + } + if (err) + abort (); + #pragma omp target map (alloc: s.u[0:5]) + { + err = 0; + if (s.u[0] != 5 || s.u[1] != 4 || s.u[2] != 5 || s.u[3] != 6 || s.u[4] != 9) + err = 1; + s.u[1] = 12; + } + #pragma omp target update from (s.s, s.u[0:5]) from (s.v[1:3]) + if (err || s.s != 11 || u[0] != 0 || u[1] != 1 || u[2] != 2 || u[3] != 5 + || u[4] != 12 || u[5] != 5 || u[6] != 6 || u[7] != 9 || u[8] != 8 + || u[9] != 9 || s.v[0] != 10 || s.v[1] != 12 || s.v[2] != 14 + || s.v[3] != 13 || s.v[4] != 14) + abort (); + #pragma omp target exit data map (release: s.s) + #pragma omp target exit data map (release: s.u[0:5]) + #pragma omp target exit data map (delete: s.v[1:3]) + #pragma omp target exit data map (release: s.s) + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-24.c b/libgomp/testsuite/libgomp.c/target-24.c new file mode 100644 index 0000000..e0ff29a --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-24.c @@ -0,0 +1,43 @@ +#include <omp.h> +#include <stdlib.h> + +int +main () +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + int a[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }; + int *b = a; + int shared_mem = 0; + #pragma omp target map (alloc: shared_mem) + shared_mem = 1; + if (omp_target_is_present (b, d) != shared_mem) + abort (); + #pragma omp target enter data map (to: a) + if (omp_target_is_present (b, d) == 0) + abort (); + #pragma omp target enter data map (alloc: b[:0]) + if (omp_target_is_present (b, d) == 0) + abort (); + #pragma omp target exit data map (release: b[:0]) + if (omp_target_is_present (b, d) == 0) + abort (); + #pragma omp target exit data map (release: b[:0]) + if (omp_target_is_present (b, d) != shared_mem) + abort (); + #pragma omp target enter data map (to: a) + if (omp_target_is_present (b, d) == 0) + abort (); + #pragma omp target enter data map (always, to: b[:0]) + if (omp_target_is_present (b, d) == 0) + abort (); + #pragma omp target exit data map (delete: b[:0]) + if (omp_target_is_present (b, d) != shared_mem) + abort (); + #pragma omp target exit data map (from: b[:0]) + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-25.c b/libgomp/testsuite/libgomp.c/target-25.c new file mode 100644 index 0000000..aeb19ae --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-25.c @@ -0,0 +1,84 @@ +#include <stdlib.h> +#include <unistd.h> + +int +main () +{ + int x = 0, y = 0, z = 0, s = 11, t = 12, u = 13, w = 7, err; + #pragma omp parallel + #pragma omp single + { + #pragma omp task depend(in: x) + { + usleep (5000); + x = 1; + } + #pragma omp task depend(in: x) + { + usleep (6000); + y = 2; + } + #pragma omp task depend(out: z) + { + usleep (7000); + z = 3; + } + #pragma omp target map(tofrom: x) map(from: err) firstprivate (y) depend(inout: x, z) + err = (x != 1 || y != 2 || z != 3); + if (err) + abort (); + #pragma omp task depend(in: x) + { + usleep (5000); + x = 4; + } + #pragma omp task depend(in: x) + { + usleep (4000); + y = 5; + } + #pragma omp task depend(in: z) + { + usleep (3000); + z = 6; + } + #pragma omp target enter data nowait map (to: w) + #pragma omp target enter data depend (inout: x, z) map (to: x, y, z) + #pragma omp target map (alloc: x, y, z) map(from: err) + { + err = (x != 4 || y != 5 || z != 6); + x = 7; + y = 8; + z = 9; + } + if (err) + abort (); + #pragma omp taskwait + #pragma omp target map (alloc: w) map(from: err) + { + err = w != 7; + w = 17; + } + if (err) + abort (); + #pragma omp task depend(in: x) + { + usleep (2000); + s = 14; + } + #pragma omp task depend(in: x) + { + usleep (3000); + t = 15; + } + #pragma omp task depend(in: z) + { + usleep (4000); + u = 16; + } + #pragma omp target exit data depend (inout: x, z) map (from: x, y, z, w) + if (x != 7 || y != 8 || z != 9 || s != 14 || t != 15 || u != 16 || w != 17) + abort (); + } + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-26.c b/libgomp/testsuite/libgomp.c/target-26.c new file mode 100644 index 0000000..fa6b525 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-26.c @@ -0,0 +1,36 @@ +extern void abort (void); +#pragma omp declare target +int a[4] = { 2, 3, 4, 5 }, *b; +#pragma omp end declare target + +int +main () +{ + int err; + int c[3] = { 6, 7, 8 }; + b = c; + #pragma omp target map(to: a[0:2], b[0:2]) map(from: err) + err = a[0] != 2 || a[1] != 3 || a[2] != 4 || a[3] != 5 || b[0] != 6 || b[1] != 7; + if (err) + abort (); + a[1] = 9; + a[2] = 10; + #pragma omp target map(always,to:a[1:2]) map(from: err) + err = a[0] != 2 || a[1] != 9 || a[2] != 10 || a[3] != 5; + if (err) + abort (); + #pragma omp parallel firstprivate(a, b, c, err) num_threads (2) + #pragma omp single + { + b = c + 1; + a[0] = 11; + a[2] = 13; + c[1] = 14; + int d = 0; + #pragma omp target map(to: a[0:3], b[d:2]) map (from: err) + err = a[0] != 11 || a[1] != 9 || a[2] != 13 || b[0] != 14 || b[1] != 8; + if (err) + abort (); + } + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-27.c b/libgomp/testsuite/libgomp.c/target-27.c new file mode 100644 index 0000000..c86651b --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-27.c @@ -0,0 +1,67 @@ +#include <stdlib.h> +#include <unistd.h> + +int +main () +{ + int x = 0, y = 0, z = 0, err; + int shared_mem = 0; + #pragma omp target map(to: shared_mem) + shared_mem = 1; + #pragma omp parallel + #pragma omp single + { + #pragma omp task depend(in: x) + { + usleep (5000); + x = 1; + } + #pragma omp task depend(in: x) + { + usleep (6000); + y = 2; + } + #pragma omp task depend(out: z) + { + usleep (7000); + z = 3; + } + #pragma omp target enter data map(to: x, y, z) depend(inout: x, z) nowait + #pragma omp task depend(inout: x, z) + { + x++; y++; z++; + } + #pragma omp target update to(x, y) depend(inout: x) nowait + #pragma omp target enter data map(always, to: z) depend(inout: z) nowait + #pragma omp target map (alloc: x, y, z) map (from: err) depend(inout: x, z) + { + err = x != 2 || y != 3 || z != 4; + x = 5; y = 6; z = 7; + } + #pragma omp task depend(in: x) + { + usleep (5000); + if (!shared_mem) + x = 1; + } + #pragma omp task depend(in: x) + { + usleep (6000); + if (!shared_mem) + y = 2; + } + #pragma omp task depend(out: z) + { + usleep (3000); + if (!shared_mem) + z = 3; + } + #pragma omp target exit data map(release: z) depend(inout: z) nowait + #pragma omp target exit data map(from: x, y) depend(inout: x) nowait + #pragma omp target exit data map(from: z) depend(inout: z) nowait + #pragma omp taskwait + if (err || x != 5 || y != 6 || z != 7) + abort (); + } + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-7.c b/libgomp/testsuite/libgomp.c/target-7.c index 0fe6150..41a1332 100644 --- a/libgomp/testsuite/libgomp.c/target-7.c +++ b/libgomp/testsuite/libgomp.c/target-7.c @@ -37,63 +37,63 @@ foo (int f) abort (); #pragma omp target data device (d) map (to: h) { - #pragma omp target device (d) + #pragma omp target device (d) map (h) if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 5) abort (); #pragma omp target update device (d) from (h) } #pragma omp target data if (v > 1) map (to: h) { - #pragma omp target if (v > 1) + #pragma omp target if (v > 1) map(h) if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 6) abort (); #pragma omp target update if (v > 1) from (h) } #pragma omp target data device (d) if (v > 1) map (to: h) { - #pragma omp target device (d) if (v > 1) + #pragma omp target device (d) if (v > 1) map(h) if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 7) abort (); #pragma omp target update device (d) if (v > 1) from (h) } #pragma omp target data if (v <= 1) map (to: h) { - #pragma omp target if (v <= 1) + #pragma omp target if (v <= 1) map (tofrom: h) if (omp_get_level () != 0 || h++ != 8) abort (); #pragma omp target update if (v <= 1) from (h) } #pragma omp target data device (d) if (v <= 1) map (to: h) { - #pragma omp target device (d) if (v <= 1) + #pragma omp target device (d) if (v <= 1) map (h) if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 9) abort (); #pragma omp target update device (d) if (v <= 1) from (h) } #pragma omp target data if (0) map (to: h) { - #pragma omp target if (0) + #pragma omp target if (0) map (h) if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 10) abort (); #pragma omp target update if (0) from (h) } #pragma omp target data device (d) if (0) map (to: h) { - #pragma omp target device (d) if (0) + #pragma omp target device (d) if (0) map (h) if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 11) abort (); #pragma omp target update device (d) if (0) from (h) } #pragma omp target data if (1) map (to: h) { - #pragma omp target if (1) + #pragma omp target if (1) map (tofrom: h) if (omp_get_level () != 0 || h++ != 12) abort (); #pragma omp target update if (1) from (h) } #pragma omp target data device (d) if (1) map (to: h) { - #pragma omp target device (d) if (1) + #pragma omp target device (d) if (1) map (tofrom: h) if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 13) abort (); #pragma omp target update device (d) if (1) from (h) diff --git a/libgomp/testsuite/libgomp.c/taskloop-1.c b/libgomp/testsuite/libgomp.c/taskloop-1.c new file mode 100644 index 0000000..21551f2 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/taskloop-1.c @@ -0,0 +1,46 @@ +/* { dg-do run } */ +/* { dg-options "-O2 -fopenmp -std=c99" } */ + +int q, r, e; + +__attribute__((noinline, noclone)) void +foo (long a, long b) +{ + #pragma omp taskloop lastprivate (q) nogroup + for (long d = a; d < b; d += 2) + { + q = d; + if (d < 2 || d > 6 || (d & 1)) + #pragma omp atomic + e |= 1; + } +} + +__attribute__((noinline, noclone)) int +bar (int a, int b) +{ + int q = 7; + #pragma omp taskloop lastprivate (q) + for (int d = a; d < b; d++) + { + if (d < 12 || d > 17) + #pragma omp atomic + e |= 1; + q = d; + } + return q; +} + +int +main () +{ + #pragma omp parallel + #pragma omp single + { + foo (2, 7); + r = bar (12, 18); + } + if (q != 6 || r != 17 || e) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/taskloop-2.c b/libgomp/testsuite/libgomp.c/taskloop-2.c new file mode 100644 index 0000000..be893eb --- /dev/null +++ b/libgomp/testsuite/libgomp.c/taskloop-2.c @@ -0,0 +1,147 @@ +/* { dg-do run } */ +/* { dg-options "-O2 -std=c99" } */ +/* { dg-additional-options "-msse2" { target sse2_runtime } } */ +/* { dg-additional-options "-mavx" { target avx_runtime } } */ + +int u[1024], v[1024], w[1024], m; + +__attribute__((noinline, noclone)) void +f1 (long a, long b) +{ + #pragma omp taskloop simd default(none) shared(u, v, w) nogroup + for (long d = a; d < b; d++) + u[d] = v[d] + w[d]; +} + +__attribute__((noinline, noclone)) int +f2 (long a, long b, long c) +{ + int d, e; + #pragma omp taskloop simd default(none) shared(u, v, w) linear(d:1) linear(c:5) lastprivate(e) + for (d = a; d < b; d++) + { + u[d] = v[d] + w[d]; + c = c + 5; + e = c + 9; + } + return d + c + e; +} + +__attribute__((noinline, noclone)) int +f3 (long a, long b) +{ + int d; + #pragma omp taskloop simd default(none) shared(u, v, w) + for (d = a; d < b; d++) + { + int *p = &d; + u[d] = v[d] + w[d]; + } + return d; +} + +__attribute__((noinline, noclone)) int +f4 (long a, long b, long c, long d) +{ + int e, f, g; + #pragma omp taskloop simd default(none) shared(u, v, w) collapse(2) lastprivate(g) + for (e = a; e < b; e++) + for (f = c; f < d; f++) + { + int *p = &e; + int *q = &f; + int r = 32 * e + f; + u[r] = v[r] + w[r]; + g = r; + } + return e + f + g; +} + +__attribute__((noinline, noclone)) int +f5 (long a, long b, long c, long d) +{ + int e, f; + #pragma omp taskloop simd default(none) shared(u, v, w) collapse(2) + for (e = a; e < b; e++) + for (f = c; f < d; f++) + { + int r = 32 * e + f; + u[r] = v[r] + w[r]; + } + return e + f; +} + +int +main () +{ + int i; + for (i = 0; i < 1024; i++) + { + v[i] = i; + w[i] = i + 1; + } + #pragma omp parallel + #pragma omp single + f1 (0, 1024); + for (i = 0; i < 1024; i++) + if (u[i] != 2 * i + 1) + __builtin_abort (); + else + { + v[i] = 1024 - i; + w[i] = 512 - i; + } + #pragma omp parallel + #pragma omp single + m = f2 (2, 1022, 17); + for (i = 0; i < 1024; i++) + if ((i < 2 || i >= 1022) ? u[i] != 2 * i + 1 : u[i] != 1536 - 2 * i) + __builtin_abort (); + else + { + v[i] = i; + w[i] = i + 1; + } + if (m != 1022 + 2 * (1020 * 5 + 17) + 9) + __builtin_abort (); + #pragma omp parallel + #pragma omp single + m = f3 (0, 1024); + for (i = 0; i < 1024; i++) + if (u[i] != 2 * i + 1) + __builtin_abort (); + else + { + v[i] = 1024 - i; + w[i] = 512 - i; + } + if (m != 1024) + __builtin_abort (); + #pragma omp parallel + #pragma omp single + m = f4 (0, 32, 0, 32); + for (i = 0; i < 1024; i++) + if (u[i] != 1536 - 2 * i) + __builtin_abort (); + else + { + v[i] = i; + w[i] = i + 1; + } + if (m != 32 + 32 + 1023) + __builtin_abort (); + #pragma omp parallel + #pragma omp single + m = f5 (0, 32, 0, 32); + for (i = 0; i < 1024; i++) + if (u[i] != 2 * i + 1) + __builtin_abort (); + else + { + v[i] = 1024 - i; + w[i] = 512 - i; + } + if (m != 32 + 32) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/taskloop-3.c b/libgomp/testsuite/libgomp.c/taskloop-3.c new file mode 100644 index 0000000..5356d7f --- /dev/null +++ b/libgomp/testsuite/libgomp.c/taskloop-3.c @@ -0,0 +1,84 @@ +/* { dg-do run } */ +/* { dg-options "-O2 -fopenmp -std=c99" } */ + +int g; +int a[1024]; + +__attribute__((noinline, noclone)) int +f1 (int x) +{ + #pragma omp taskloop firstprivate (x) lastprivate (x) + for (int i = 0; i < 64; i++) + { + if (x != 74) + __builtin_abort (); + if (i == 63) + x = i + 4; + } + return x; +} + +__attribute__((noinline, noclone)) void +f2 (void) +{ + #pragma omp taskloop firstprivate (g) lastprivate (g) nogroup + for (int i = 0; i < 64; i++) + { + if (g != 77) + __builtin_abort (); + if (i == 63) + g = i + 9; + } +} + +__attribute__((noinline, noclone)) long long +f3 (long long a, long long b, long long c) +{ + long long i; + int l; + #pragma omp taskloop default (none) lastprivate (i, l) + for (i = a; i < b; i += c) + l = i; + return l * 7 + i; +} + +__attribute__((noinline, noclone)) long long +f4 (long long a, long long b, long long c, long long d, + long long e, long long f, int k) +{ + long long i, j; + int l; + #pragma omp taskloop default (none) collapse(2) \ + firstprivate (k) lastprivate (i, j, k, l) + for (i = a; i < b; i += e) + for (j = c; j < d; j += f) + { + if (k != 73) + __builtin_abort (); + if (i == 31 && j == 46) + k = i; + l = j; + } + return i + 5 * j + 11 * k + 17 * l; +} + +int +main () +{ + #pragma omp parallel + #pragma omp single + { + if (f1 (74) != 63 + 4) + __builtin_abort (); + g = 77; + f2 (); + #pragma omp taskwait + if (g != 63 + 9) + __builtin_abort (); + if (f3 (7, 12, 2) != 11 * 7 + 13) + __builtin_abort (); + if (f4 (0, 32, 16, 48, 1, 2, 73) != 32 + 5 * 48 + 11 * 31 + 17 * 46) + __builtin_abort (); + } + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/taskloop-4.c b/libgomp/testsuite/libgomp.c/taskloop-4.c new file mode 100644 index 0000000..a69be19 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/taskloop-4.c @@ -0,0 +1,97 @@ +/* { dg-do run } */ +/* { dg-options "-O2 -fopenmp" } */ + +int u[64], v; + +__attribute__((noinline, noclone)) int +test (int a, int b, int c, int d, void (*fn) (int, int, int, int), + int *num_tasks, int *min_iters, int *max_iters) +{ + int i, t = 0; + __builtin_memset (u, 0, sizeof u); + v = 0; + fn (a, b, c, d); + *min_iters = 0; + *max_iters = 0; + *num_tasks = v; + if (v) + { + *min_iters = u[0]; + *max_iters = u[0]; + t = u[0]; + for (i = 1; i < v; i++) + { + if (*min_iters > u[i]) + *min_iters = u[i]; + if (*max_iters < u[i]) + *max_iters = u[i]; + t += u[i]; + } + } + return t; +} + +void +grainsize (int a, int b, int c, int d) +{ + int i, j = 0, k = 0; + #pragma omp taskloop firstprivate (j, k) grainsize(d) + for (i = a; i < b; i += c) + { + if (j == 0) + { + #pragma omp atomic capture + k = v++; + if (k >= 64) + __builtin_abort (); + } + u[k] = ++j; + } +} + +void +num_tasks (int a, int b, int c, int d) +{ + int i, j = 0, k = 0; + #pragma omp taskloop firstprivate (j, k) num_tasks(d) + for (i = a; i < b; i += c) + { + if (j == 0) + { + #pragma omp atomic capture + k = v++; + if (k >= 64) + __builtin_abort (); + } + u[k] = ++j; + } +} + +int +main () +{ + #pragma omp parallel + #pragma omp single + { + int min_iters, max_iters, ntasks; + /* If grainsize is present, # of task loop iters is >= grainsize && < 2 * grainsize, + unless # of loop iterations is smaller than grainsize. */ + if (test (0, 79, 1, 17, grainsize, &ntasks, &min_iters, &max_iters) != 79 + || min_iters < 17 || max_iters >= 17 * 2) + __builtin_abort (); + if (test (-49, 2541, 7, 28, grainsize, &ntasks, &min_iters, &max_iters) != 370 + || min_iters < 28 || max_iters >= 28 * 2) + __builtin_abort (); + if (test (7, 21, 2, 15, grainsize, &ntasks, &min_iters, &max_iters) != 7 + || ntasks != 1 || min_iters != 7 || max_iters != 7) + __builtin_abort (); + /* If num_tasks is present, # of task loop iters is min (# of loop iters, num_tasks). */ + if (test (-51, 2500, 48, 9, num_tasks, &ntasks, &min_iters, &max_iters) != 54 + || ntasks != 9) + __builtin_abort (); + if (test (0, 25, 2, 17, num_tasks, &ntasks, &min_iters, &max_iters) != 13 + || ntasks != 13) + __builtin_abort (); + } + return 0; +} |