diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 9e393803eb2..7f7fc13e139 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,9 @@ +2016-07-15 Cesar Philippidis + + * omp-low.c (lower_omp_target): Mark data clauses with + GOMP_MAP_FORCE_{PRESENT,TO,FROM,TOFROM} as potentially having + zero-length subarrays. + 2016-07-15 Richard Biener PR tree-optimization/71881 diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index 968d942c7be..bc1e9558774 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,11 @@ +2016-07-15 Cesar Philippidis + + * 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. + 2016-07-15 Jakub Jelinek PR c/71858 diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 8fb4e4dd14f..ef585fb9990 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -13664,11 +13664,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 bafd0d23195..0b5ab91d0d5 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -12505,6 +12505,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/ChangeLog b/gcc/cp/ChangeLog index f7f7ba3ce34..f85a3333281 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,11 @@ +2016-07-15 Cesar Philippidis + + * 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. + 2016-07-11 Jason Merrill * decl.c (store_parm_decls): Remove check for void parm. diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index ef35aa9cfb0..3e865b0fb64 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -35258,11 +35258,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 d1fb11973fd..615d3aeab2a 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5002,6 +5002,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 ecba096c5c1..04509f39938 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -16210,6 +16210,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/ChangeLog b/libgomp/ChangeLog index 35e29c06776..e4555eacd2d 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,8 @@ +2016-07-15 Cesar Philippidis + + * testsuite/libgomp.oacc-c-c++-common/zero_length_subarrays.c: New + test. + 2016-07-03 H.J. Lu PR middle-end/71734 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 00000000000..8954551205b --- /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; +}