[PATCH] zero-length arrays in OpenACC

[Cesar Philippidis <cesar@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 402 bytes --]

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

[-- Attachment #2: subarray-followup.diff --]
[-- Type: text/x-patch, Size: 4201 bytes --]

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;
+}