diff mbox

zero-length arrays in OpenACC

Message ID 574F5530.9080408@codesourcery.com
State New
Headers show

Commit Message

Cesar Philippidis June 1, 2016, 9:35 p.m. UTC 
This patch teaches c and c++ front ends and omp-low how to deal with
subarray involving GOMP_MAP_FORCE_{PRESENT,TO,FROM,TOFROM} data
mappings. As the libgomp test case shows, it might be possible for a
subarray to have zero length. This patch fixes that.

I also updated *parser_oacc_declare not to handle GOMP_MAP_POINTER,
because subarray pointers are now firstprivate.

Is this OK for trunk?

Cesar

Comments

Cesar Philippidis June 16, 2016, 2:24 a.m. UTC | #1 
Ping.

Cesar

On 06/01/2016 02:35 PM, Cesar Philippidis wrote:
> This patch teaches c and c++ front ends and omp-low how to deal with
> subarray involving GOMP_MAP_FORCE_{PRESENT,TO,FROM,TOFROM} data
> mappings. As the libgomp test case shows, it might be possible for a
> subarray to have zero length. This patch fixes that.
> 
> I also updated *parser_oacc_declare not to handle GOMP_MAP_POINTER,
> because subarray pointers are now firstprivate.
> 
> Is this OK for trunk?
> 
> Cesar
>
Cesar Philippidis July 15, 2016, 1:57 a.m. UTC | #2 
Ping, x2.

Cesar

On 06/01/2016 02:35 PM, Cesar Philippidis wrote:
> This patch teaches c and c++ front ends and omp-low how to deal with
> subarray involving GOMP_MAP_FORCE_{PRESENT,TO,FROM,TOFROM} data
> mappings. As the libgomp test case shows, it might be possible for a
> subarray to have zero length. This patch fixes that.
> 
> I also updated *parser_oacc_declare not to handle GOMP_MAP_POINTER,
> because subarray pointers are now firstprivate.
> 
> Is this OK for trunk?
> 
> Cesar
>
Jakub Jelinek July 15, 2016, 8:35 a.m. UTC | #3 
On Wed, Jun 01, 2016 at 02:35:44PM -0700, Cesar Philippidis wrote:
> This patch teaches c and c++ front ends and omp-low how to deal with
> subarray involving GOMP_MAP_FORCE_{PRESENT,TO,FROM,TOFROM} data
> mappings. As the libgomp test case shows, it might be possible for a
> subarray to have zero length. This patch fixes that.
> 
> I also updated *parser_oacc_declare not to handle GOMP_MAP_POINTER,
> because subarray pointers are now firstprivate.
> 
> Is this OK for trunk?

Ok with me if that is the semantics OpenACC wants for zero length array
sections (which would surprise me, but it is up to you to discuss the
semantics you want).
Basically, in OpenMP 4.5, it has been voted in that zero length array
sections don't cause mapping, but rather just resolve its address to
either some device pointer if the array section points into an already
mapped object, or NULL if it corresponds to something not mapped yet.
As GNU extension, GCC still handles zero length objects (not array sections)
differently (by mapping them).

	Jakub
diff mbox

Patch

2016-06-01  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/c/
	* c-parser.c (c_parser_oacc_declare): Don't scan for
	GOMP_MAP_POINTER.
	* c-typeck.c (handle_omp_array_sections): Mark data clauses with
	GOMP_MAP_FORCE_{PRESENT,TO,FROM,TOFROM} as potentially having
	zero-length subarrays.

	gcc/cp/
	* parser.c (cp_parser_oacc_declare): Don't scan for
	GOMP_MAP_POINTER.
	* semantics.c (handle_omp_array_sections): Mark data clauses with
	GOMP_MAP_FORCE_{PRESENT,TO,FROM,TOFROM} as potentially having
	zero-length subarrays.

	gcc/
	* omp-low.c (lower_omp_target): Mark data clauses with
	GOMP_MAP_FORCE_{PRESENT,TO,FROM,TOFROM} as potentially having
	zero-length subarrays.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/zero_length_subarrays.c: New
	test.


diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index bca8653..9bd377c 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -13622,11 +13622,6 @@  c_parser_oacc_declare (c_parser *parser)
 	case GOMP_MAP_DEVICE_RESIDENT:
 	  break;
 
-	case GOMP_MAP_POINTER:
-	  /* Generated by c_finish_omp_clauses from array sections;
-	     avoid spurious diagnostics.  */
-	  break;
-
 	case GOMP_MAP_LINK:
 	  if (!global_bindings_p ()
 	      && (TREE_STATIC (decl)
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 1520c20..1e97749 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -12439,6 +12439,10 @@  handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	  case GOMP_MAP_ALWAYS_TOFROM:
 	  case GOMP_MAP_RELEASE:
 	  case GOMP_MAP_DELETE:
+	  case GOMP_MAP_FORCE_TO:
+	  case GOMP_MAP_FORCE_FROM:
+	  case GOMP_MAP_FORCE_TOFROM:
+	  case GOMP_MAP_FORCE_PRESENT:
 	    OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
 	    break;
 	  default:
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 29a1b80..9568ef8 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -35244,11 +35244,6 @@  cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
 	case GOMP_MAP_DEVICE_RESIDENT:
 	  break;
 
-	case GOMP_MAP_POINTER:
-	  /* Generated by c_finish_omp_clauses from array sections;
-	     avoid spurious diagnostics.  */
-	  break;
-
 	case GOMP_MAP_LINK:
 	  if (!global_bindings_p ()
 	      && (TREE_STATIC (decl)
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 8e682c5..d960747 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -4995,6 +4995,10 @@  handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	      case GOMP_MAP_ALWAYS_TOFROM:
 	      case GOMP_MAP_RELEASE:
 	      case GOMP_MAP_DELETE:
+	      case GOMP_MAP_FORCE_TO:
+	      case GOMP_MAP_FORCE_FROM:
+	      case GOMP_MAP_FORCE_TOFROM:
+	      case GOMP_MAP_FORCE_PRESENT:
 		OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
 		break;
 	      default:
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 77bdb18..052e6d3 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -16208,6 +16208,10 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		    case GOMP_MAP_ALWAYS_FROM:
 		    case GOMP_MAP_ALWAYS_TOFROM:
 		    case GOMP_MAP_RELEASE:
+		    case GOMP_MAP_FORCE_TO:
+		    case GOMP_MAP_FORCE_FROM:
+		    case GOMP_MAP_FORCE_TOFROM:
+		    case GOMP_MAP_FORCE_PRESENT:
 		      tkind_zero = GOMP_MAP_ZERO_LEN_ARRAY_SECTION;
 		      break;
 		    case GOMP_MAP_DELETE:
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/zero_length_subarrays.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/zero_length_subarrays.c
new file mode 100644
index 0000000..8954551
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/zero_length_subarrays.c
@@ -0,0 +1,45 @@ 
+/* Exercise zero-length sub-arrays.  */
+
+const int n = 10;
+
+void
+subzero_present (int *a, int n)
+{
+#pragma acc data present (a[0:n])
+  ;
+#pragma acc data pcopy (a[0:n])
+  ;
+#pragma acc data pcopyin (a[0:n])
+  ;
+#pragma acc data pcopyout (a[0:n])
+  ;
+
+}
+
+void
+subzero (int *a, int n)
+{
+#pragma acc data create (a[0:n])
+  ;
+#pragma acc data copy (a[0:n])
+  ;
+#pragma acc data copyin (a[0:n])
+  ;
+#pragma acc data copyout (a[0:n])
+  ;
+}
+
+int
+main ()
+{
+  int a[n];
+
+#pragma acc data copy (a[0:n])
+  {
+    subzero_present (a, 0);
+  }
+
+  subzero (a, 0);
+
+  return 0;
+}