[gcc/devel/omp/gcc-9] [8/8] Multi-dimensional dynamic array support for OpenACC data clauses, libgomp testsuite additions
Thomas Schwinge
tschwinge@gcc.gnu.org
Tue Mar 3 12:03:00 GMT 2020
https://gcc.gnu.org/g:c42632498cc6a3aeb0cbb4375ede777b0bcc1d38
commit c42632498cc6a3aeb0cbb4375ede777b0bcc1d38
Author: Julian Brown <julian@codesourcery.com>
Date: Tue Feb 5 12:51:36 2019 -0800
[8/8] Multi-dimensional dynamic array support for OpenACC data clauses, libgomp testsuite additions
2018-10-16 Chung-Lin Tang <cltang@codesourcery.com>
libgomp/
* testsuite/libgomp.oacc-c-c++-common/da-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/da-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/da-3.c: New test.
* testsuite/libgomp.oacc-c-c++-common/da-4.c: New test.
* testsuite/libgomp.oacc-c-c++-common/da-utils.h: New test.
(cherry picked from openacc-gcc-9-branch commit
26413f817be2e0e0ec02333930c3d01acea4f1c7)
Diff:
---
libgomp/ChangeLog.omp | 8 ++
libgomp/testsuite/libgomp.oacc-c-c++-common/da-1.c | 103 +++++++++++++++++++++
libgomp/testsuite/libgomp.oacc-c-c++-common/da-2.c | 37 ++++++++
libgomp/testsuite/libgomp.oacc-c-c++-common/da-3.c | 45 +++++++++
libgomp/testsuite/libgomp.oacc-c-c++-common/da-4.c | 36 +++++++
.../testsuite/libgomp.oacc-c-c++-common/da-utils.h | 44 +++++++++
6 files changed, 273 insertions(+)
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index d1df043..990199a 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,5 +1,13 @@
2018-10-16 Chung-Lin Tang <cltang@codesourcery.com>
+ * testsuite/libgomp.oacc-c-c++-common/da-1.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/da-2.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/da-3.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/da-4.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/da-utils.h: New test.
+
+2018-10-16 Chung-Lin Tang <cltang@codesourcery.com>
+
* target.c (struct da_dim): New struct declaration.
(struct da_descr_type): Likewise.
(struct da_info): Likewise.
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/da-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-1.c
new file mode 100644
index 0000000..c1c205d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-1.c
@@ -0,0 +1,103 @@
+/* { dg-do run { target { ! openacc_host_selected } } } */
+
+#include <stdlib.h>
+#include <assert.h>
+
+#define n 100
+#define m 100
+
+int b[n][m];
+
+void
+test1 (void)
+{
+ int i, j, *a[100];
+
+ /* Array of pointers form test. */
+ for (i = 0; i < n; i++)
+ {
+ a[i] = (int *)malloc (sizeof (int) * m);
+ for (j = 0; j < m; j++)
+ b[i][j] = j - i;
+ }
+
+ #pragma acc parallel loop copyout(a[0:n][0:m]) copyin(b)
+ for (i = 0; i < n; i++)
+ #pragma acc loop
+ for (j = 0; j < m; j++)
+ a[i][j] = b[i][j];
+
+ for (i = 0; i < n; i++)
+ {
+ for (j = 0; j < m; j++)
+ assert (a[i][j] == b[i][j]);
+ /* Clean up. */
+ free (a[i]);
+ }
+}
+
+void
+test2 (void)
+{
+ int i, j, **a = (int **) malloc (sizeof (int *) * n);
+
+ /* Separately allocated blocks. */
+ for (i = 0; i < n; i++)
+ {
+ a[i] = (int *)malloc (sizeof (int) * m);
+ for (j = 0; j < m; j++)
+ b[i][j] = j - i;
+ }
+
+ #pragma acc parallel loop copyout(a[0:n][0:m]) copyin(b)
+ for (i = 0; i < n; i++)
+ #pragma acc loop
+ for (j = 0; j < m; j++)
+ a[i][j] = b[i][j];
+
+ for (i = 0; i < n; i++)
+ {
+ for (j = 0; j < m; j++)
+ assert (a[i][j] == b[i][j]);
+ /* Clean up. */
+ free (a[i]);
+ }
+ free (a);
+}
+
+void
+test3 (void)
+{
+ int i, j, **a = (int **) malloc (sizeof (int *) * n);
+ a[0] = (int *) malloc (sizeof (int) * n * m);
+
+ /* Rows allocated in one contiguous block. */
+ for (i = 0; i < n; i++)
+ {
+ a[i] = *a + i * m;
+ for (j = 0; j < m; j++)
+ b[i][j] = j - i;
+ }
+
+ #pragma acc parallel loop copyout(a[0:n][0:m]) copyin(b)
+ for (i = 0; i < n; i++)
+ #pragma acc loop
+ for (j = 0; j < m; j++)
+ a[i][j] = b[i][j];
+
+ for (i = 0; i < n; i++)
+ for (j = 0; j < m; j++)
+ assert (a[i][j] == b[i][j]);
+
+ free (a[0]);
+ free (a);
+}
+
+int
+main (void)
+{
+ test1 ();
+ test2 ();
+ test3 ();
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/da-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-2.c
new file mode 100644
index 0000000..6ee7855
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-2.c
@@ -0,0 +1,37 @@
+/* { dg-do run { target { ! openacc_host_selected } } } */
+
+#include <assert.h>
+#include "da-utils.h"
+
+int
+main (void)
+{
+ int n = 10;
+ int ***a = (int ***) create_da (sizeof (int), n, 3);
+ int ***b = (int ***) create_da (sizeof (int), n, 3);
+ int ***c = (int ***) create_da (sizeof (int), n, 3);
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ {
+ a[i][j][k] = i + j * k + k;
+ b[i][j][k] = j + k * i + i * j;
+ c[i][j][k] = a[i][j][k];
+ }
+
+ #pragma acc parallel copy (a[0:n][0:n][0:n]) copyin (b[0:n][0:n][0:n])
+ {
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ a[i][j][k] += b[k][j][i] + i + j + k;
+ }
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ assert (a[i][j][k] == c[i][j][k] + b[k][j][i] + i + j + k);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/da-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-3.c
new file mode 100644
index 0000000..877c6df
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-3.c
@@ -0,0 +1,45 @@
+/* { dg-do run { target { ! openacc_host_selected } } } */
+
+#include <assert.h>
+#include "da-utils.h"
+
+int main (void)
+{
+ int n = 20, x = 5, y = 12;
+ int *****a = (int *****) create_da (sizeof (int), n, 5);
+
+ int sum1 = 0, sum2 = 0, sum3 = 0;
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ for (int l = 0; l < n; l++)
+ for (int m = 0; m < n; m++)
+ {
+ a[i][j][k][l][m] = 1;
+ sum1++;
+ }
+
+ #pragma acc parallel copy (a[x:y][x:y][x:y][x:y][x:y]) copy(sum2)
+ {
+ for (int i = x; i < x + y; i++)
+ for (int j = x; j < x + y; j++)
+ for (int k = x; k < x + y; k++)
+ for (int l = x; l < x + y; l++)
+ for (int m = x; m < x + y; m++)
+ {
+ a[i][j][k][l][m] = 0;
+ sum2++;
+ }
+ }
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ for (int l = 0; l < n; l++)
+ for (int m = 0; m < n; m++)
+ sum3 += a[i][j][k][l][m];
+
+ assert (sum1 == sum2 + sum3);
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/da-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-4.c
new file mode 100644
index 0000000..2059c5f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-4.c
@@ -0,0 +1,36 @@
+/* { dg-do run { target { ! openacc_host_selected } } } */
+
+#include <assert.h>
+#include "da-utils.h"
+
+int main (void)
+{
+ int n = 128;
+ double ***a = (double ***) create_da (sizeof (double), n, 3);
+ double ***b = (double ***) create_da (sizeof (double), n, 3);
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ a[i][j][k] = i + j + k + i * j * k;
+
+ /* This test exercises async copyout of dynamic array rows. */
+ #pragma acc parallel copyin(a[0:n][0:n][0:n]) copyout(b[0:n][0:n][0:n]) async(5)
+ {
+ #pragma acc loop gang
+ for (int i = 0; i < n; i++)
+ #pragma acc loop vector
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ b[i][j][k] = a[i][j][k] * 2.0;
+ }
+
+ #pragma acc wait (5)
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ assert (b[i][j][k] == a[i][j][k] * 2.0);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/da-utils.h b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-utils.h
new file mode 100644
index 0000000..2f87795
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/da-utils.h
@@ -0,0 +1,44 @@
+#include <stdlib.h>
+#include <string.h>
+#include <assert.h>
+#include <stdint.h>
+
+/* Allocate and create a pointer based NDIMS-dimensional array,
+ each dimension DIMLEN long, with ELSIZE sized data elements. */
+void *
+create_da (size_t elsize, int dimlen, int ndims)
+{
+ size_t blk_size = 0;
+ size_t n = 1;
+
+ for (int i = 0; i < ndims - 1; i++)
+ {
+ n *= dimlen;
+ blk_size += sizeof (void *) * n;
+ }
+ size_t data_rows_num = n;
+ size_t data_rows_offset = blk_size;
+ blk_size += elsize * n * dimlen;
+
+ void *blk = (void *) malloc (blk_size);
+ memset (blk, 0, blk_size);
+ void **curr_dim = (void **) blk;
+ n = 1;
+
+ for (int d = 0; d < ndims - 1; d++)
+ {
+ uintptr_t next_dim = (uintptr_t) (curr_dim + n * dimlen);
+ size_t next_dimlen = dimlen * (d < ndims - 2 ? sizeof (void *) : elsize);
+
+ for (int b = 0; b < n; b++)
+ for (int i = 0; i < dimlen; i++)
+ if (d < ndims - 1)
+ curr_dim[b * dimlen + i]
+ = (void*) (next_dim + b * dimlen * next_dimlen + i * next_dimlen);
+
+ n *= dimlen;
+ curr_dim = (void**) next_dim;
+ }
+ assert (n == data_rows_num);
+ return blk;
+}
More information about the Gcc-cvs
mailing list