public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [PATCH] zero-length arrays in OpenACC
@ 2016-06-01 21:35 Cesar Philippidis
  2016-06-16  2:24 ` [PATCH, ping] " Cesar Philippidis
                   ` (2 more replies)
  0 siblings, 3 replies; 4+ messages in thread
From: Cesar Philippidis @ 2016-06-01 21:35 UTC (permalink / raw)
  To: gcc-patches, Jakub Jelinek

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

^ permalink raw reply	[flat|nested] 4+ messages in thread
end of thread, other threads:[~2016-07-15  8:35 UTC | newest]

Thread overview: 4+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2016-06-01 21:35 [PATCH] zero-length arrays in OpenACC Cesar Philippidis
2016-06-16  2:24 ` [PATCH, ping] " Cesar Philippidis
2016-07-15  1:57 ` [PATCH] " Cesar Philippidis
2016-07-15  8:35 ` Jakub Jelinek

This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).