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.
@@ -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)
@@ -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:
@@ -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)
@@ -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:
@@ -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:
new file mode 100644
@@ -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;
+}