This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[PATCH] zero-length arrays in OpenACC
- From: Cesar Philippidis <cesar at codesourcery dot com>
- To: "gcc-patches at gcc dot gnu dot org" <gcc-patches at gcc dot gnu dot org>, Jakub Jelinek <jakub at redhat dot com>
- Date: Wed, 1 Jun 2016 14:35:44 -0700
- Subject: [PATCH] zero-length arrays in OpenACC
- Authentication-results: sourceware.org; auth=none
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
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;
+}