[gcc/devel/c++-modules] openacc: Set bias to zero for explicit attach/detach clauses in C and C++

Nathan Sidwell nathan@gcc.gnu.org
Mon Jul 13 14:58:55 GMT 2020


https://gcc.gnu.org/g:0d00fe404c162ad0cf922ca8455aa23a74042b63

commit 0d00fe404c162ad0cf922ca8455aa23a74042b63
Author: Julian Brown <julian@codesourcery.com>
Date:   Tue Jun 9 06:21:34 2020 -0700

    openacc: Set bias to zero for explicit attach/detach clauses in C and C++
    
    This is a fix for the pointer (or array) size inadvertently being used
    for the bias with attach and detach mapping kinds, for both C and C++.
    
    2020-07-09  Julian Brown  <julian@codesourcery.com>
                Thomas Schwinge  <thomas@codesourcery.com>
    
    gcc/c/
            PR middle-end/95270
            * c-typeck.c (c_finish_omp_clauses): Set OMP_CLAUSE_SIZE (bias) to zero
            for standalone attach/detach clauses.
    
    gcc/cp/
            PR middle-end/95270
            * semantics.c (finish_omp_clauses): Likewise.
    
    include/
            PR middle-end/95270
            * gomp-constants.h (gomp_map_kind): Expand comment for attach/detach
            mapping kinds.
    
    gcc/testsuite/
            PR middle-end/95270
            * c-c++-common/goacc/mdc-1.c: Update expected dump output for zero
            bias.
    
    libgomp/
            PR middle-end/95270
            * testsuite/libgomp.oacc-c-c++-common/pr95270-1.c: New test.
            * testsuite/libgomp.oacc-c-c++-common/pr95270-2.c: New test.

Diff:
---
 gcc/c/c-typeck.c                                   | 16 ++++++++
 gcc/cp/semantics.c                                 | 16 ++++++++
 gcc/testsuite/c-c++-common/goacc/mdc-1.c           | 14 +++----
 include/gomp-constants.h                           |  7 +++-
 .../libgomp.oacc-c-c++-common/pr95270-1.c          | 46 +++++++++++++++++++++
 .../libgomp.oacc-c-c++-common/pr95270-2.c          | 48 ++++++++++++++++++++++
 6 files changed, 139 insertions(+), 8 deletions(-)

diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index b28c2c5ff62..fb5c288b549 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -14579,6 +14579,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		}
 	      if (c_oacc_check_attachments (c))
 		remove = true;
+	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		  && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+		      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+		/* In this case, we have a single array element which is a
+		   pointer, and we already set OMP_CLAUSE_SIZE in
+		   handle_omp_array_sections above.  For attach/detach clauses,
+		   reset the OMP_CLAUSE_SIZE (representing a bias) to zero
+		   here.  */
+		OMP_CLAUSE_SIZE (c) = size_zero_node;
 	      break;
 	    }
 	  if (t == error_mark_node)
@@ -14592,6 +14601,13 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      remove = true;
 	      break;
 	    }
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	      && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+		  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+	    /* For attach/detach clauses, set OMP_CLAUSE_SIZE (representing a
+	       bias) to zero here, so it is not set erroneously to the pointer
+	       size later on in gimplify.c.  */
+	    OMP_CLAUSE_SIZE (c) = size_zero_node;
 	  if (TREE_CODE (t) == COMPONENT_REF
 	      && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
 	    {
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index d63cea96e23..4a3ef3d2839 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -7362,6 +7362,15 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		}
 	      if (cp_oacc_check_attachments (c))
 		remove = true;
+	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		  && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+		      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+		/* In this case, we have a single array element which is a
+		   pointer, and we already set OMP_CLAUSE_SIZE in
+		   handle_omp_array_sections above.  For attach/detach clauses,
+		   reset the OMP_CLAUSE_SIZE (representing a bias) to zero
+		   here.  */
+		OMP_CLAUSE_SIZE (c) = size_zero_node;
 	      break;
 	    }
 	  if (t == error_mark_node)
@@ -7375,6 +7384,13 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      remove = true;
 	      break;
 	    }
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	      && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+		  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+	    /* For attach/detach clauses, set OMP_CLAUSE_SIZE (representing a
+	       bias) to zero here, so it is not set erroneously to the pointer
+	       size later on in gimplify.c.  */
+	    OMP_CLAUSE_SIZE (c) = size_zero_node;
 	  if (REFERENCE_REF_P (t)
 	      && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
 	    {
diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
index fb5841a709d..337c1f7cc77 100644
--- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
@@ -45,12 +45,12 @@ t1 ()
 
 /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */
 /* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 8.. map.tofrom:s .len: 32" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
 /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .bias: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 0.." 1 "omplower" } } */
 /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.release:a .len: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .bias: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .bias: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .bias: 0.." 1 "omplower" } } */
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index b42b41403aa..7e44238ae03 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -139,7 +139,12 @@ enum gomp_map_kind
     /* Decrement usage count and deallocate if zero.  */
     GOMP_MAP_RELEASE =			(GOMP_MAP_FLAG_SPECIAL_2
 					 | GOMP_MAP_DELETE),
-    /* In OpenACC, attach a pointer to a mapped struct field.  */
+    /* The attach/detach mappings below use the OMP_CLAUSE_SIZE field as a
+       bias.  This will typically be zero, except when mapping an array slice
+       with a non-zero base.  In that case the bias will indicate the
+       (positive) difference between the start of the actual mapped data and
+       the "virtual" origin of the array.
+       In OpenACC, attach a pointer to a mapped struct field.  */
     GOMP_MAP_ATTACH =			(GOMP_MAP_DEEP_COPY | 0),
     /* In OpenACC, detach a pointer to a mapped struct field.  */
     GOMP_MAP_DETACH =			(GOMP_MAP_DEEP_COPY | 1),
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-1.c
new file mode 100644
index 00000000000..0457c232bc9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-1.c
@@ -0,0 +1,46 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <assert.h>
+#include <openacc.h>
+#include <stdint.h>
+
+int
+main ()
+{
+  int data;
+  int *data_p_dev = (int *) acc_create (&data, sizeof data);
+  int *data_p = &data;
+  uintptr_t ptrbits;
+
+  acc_copyin (&data_p, sizeof data_p);
+
+  /* Test attach/detach directives.  */
+#pragma acc enter data attach(data_p)
+#pragma acc serial copyout(ptrbits) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */
+  {
+    ptrbits = (uintptr_t) data_p;
+  }
+#pragma acc exit data detach(data_p)
+  assert ((void *) ptrbits == data_p_dev);
+
+  acc_update_self (&data_p, sizeof data_p);
+  assert (data_p == &data);
+
+  /* Test attach/detach API call.  */
+  acc_attach ((void **) &data_p);
+#pragma acc serial copyout(ptrbits) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */
+  {
+    ptrbits = (uintptr_t) data_p;
+  }
+  acc_detach ((void **) &data_p);
+
+  assert ((void *) ptrbits == data_p_dev);
+  acc_update_self (&data_p, sizeof data_p);
+  assert (data_p == &data);
+
+  acc_delete (&data_p, sizeof data_p);
+  acc_delete (&data, sizeof data);
+
+  return 0;
+}
+
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-2.c
new file mode 100644
index 00000000000..0575e726738
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-2.c
@@ -0,0 +1,48 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <assert.h>
+#include <openacc.h>
+#include <stdint.h>
+
+#define N 128
+
+int
+main ()
+{
+  int *ptrarr[N];
+  int otherarr[N];
+  int sum = 0, hostsum = 0;
+
+  for (int i = 0; i < N; i++)
+    {
+      otherarr[i] = i * 2 + 1;
+      ptrarr[i] = &otherarr[N - 1 - i];
+      hostsum += otherarr[i];
+    }
+
+  acc_copyin (otherarr, sizeof otherarr);
+  acc_copyin (ptrarr, sizeof ptrarr);
+
+  for (int i = 0; i < N; i++)
+    {
+      #pragma acc enter data attach(ptrarr[i])
+    }
+
+  #pragma acc parallel loop copyin(ptrarr[0:N], otherarr[0:N]) \
+		       reduction(+:sum)
+  for (int i = 0; i < N; i++)
+    sum += *ptrarr[i];
+
+  for (int i = 0; i < N; i++)
+    {
+      #pragma acc exit data detach(ptrarr[i])
+    }
+
+  assert (sum == hostsum);
+
+  acc_delete (ptrarr, sizeof ptrarr);
+  acc_delete (otherarr, sizeof otherarr);
+
+  return 0;
+}
+


More information about the Gcc-cvs mailing list