2016-06-01 Cesar Philippidis 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; +}