On Fri, Sep 04, 2015 at 09:07:02PM +0300, Ilya Verbin wrote: > It seems that there is a bug some here: > > Here is the reproducer: > > #pragma omp declare target > int a[1]; > #pragma omp end declare target > > void foo () > { > #pragma omp target map(to: a[0:1]) > a; > } > > > lookup_decl (var, ctx) tries to lookup for 'a', but ctx->cb.decl_map->get () > returns null-pointer.
Fixed thusly, tested on x86_64-linux, committed to gomp4.1 branch. 2015-09-07 Jakub Jelinek <ja...@redhat.com> * omp-low.c (scan_sharing_clauses): Don't ignore map with declare target vars for GOMP_MAP_FIRSTPRIVATE_POINTER, unless the decl is an array. (lower_omp_target): Ignore GOMP_MAP_FIRSTPRIVATE_POINTER map with declare target var if it is an array. * testsuite/libgomp.c/target-26.c: New test. --- gcc/omp-low.c.jj 2015-09-04 11:34:45.000000000 +0200 +++ gcc/omp-low.c 2015-09-07 14:10:36.198517500 +0200 @@ -2060,6 +2060,8 @@ scan_sharing_clauses (tree clauses, omp_ directly. */ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && DECL_P (decl) + && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER + || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) && varpool_node::get_create (decl)->offloadable) break; @@ -2284,6 +2286,8 @@ scan_sharing_clauses (tree clauses, omp_ break; decl = OMP_CLAUSE_DECL (c); if (DECL_P (decl) + && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER + || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) && varpool_node::get_create (decl)->offloadable) break; @@ -13358,6 +13362,10 @@ lower_omp_target (gimple_stmt_iterator * { if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) { + if (is_global_var (maybe_lookup_decl_in_outer_ctx (var, ctx)) + && varpool_node::get_create (var)->offloadable) + continue; + tree type = build_pointer_type (TREE_TYPE (var)); tree new_var = lookup_decl (var, ctx); x = create_tmp_var_raw (type, get_name (new_var)); @@ -14081,6 +14089,12 @@ lower_omp_target (gimple_stmt_iterator * HOST_WIDE_INT offset = 0; gcc_assert (prev); var = OMP_CLAUSE_DECL (c); + if (DECL_P (var) + && TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE + && is_global_var (maybe_lookup_decl_in_outer_ctx (var, + ctx)) + && varpool_node::get_create (var)->offloadable) + break; if (TREE_CODE (var) == INDIRECT_REF && TREE_CODE (TREE_OPERAND (var, 0)) == COMPONENT_REF) var = TREE_OPERAND (var, 0); --- libgomp/testsuite/libgomp.c/target-26.c.jj 2015-09-07 11:59:08.665425993 +0200 +++ libgomp/testsuite/libgomp.c/target-26.c 2015-09-07 12:38:56.000000000 +0200 @@ -0,0 +1,36 @@ +extern void abort (void); +#pragma omp declare target +int a[4] = { 2, 3, 4, 5 }, *b; +#pragma omp end declare target + +int +main () +{ + int err; + int c[3] = { 6, 7, 8 }; + b = c; + #pragma omp target map(to: a[0:2], b[0:2]) map(from: err) + err = a[0] != 2 || a[1] != 3 || a[2] != 4 || a[3] != 5 || b[0] != 6 || b[1] != 7; + if (err) + abort (); + a[1] = 9; + a[2] = 10; + #pragma omp target map(always,to:a[1:2]) map(from: err) + err = a[0] != 2 || a[1] != 9 || a[2] != 10 || a[3] != 5; + if (err) + abort (); + #pragma omp parallel firstprivate(a, b, c, err) num_threads (2) + #pragma omp single + { + b = c + 1; + a[0] = 11; + a[2] = 13; + c[1] = 14; + int d = 0; + #pragma omp target map(to: a[0:3], b[d:2]) map (from: err) + err = a[0] != 11 || a[1] != 9 || a[2] != 13 || b[0] != 14 || b[1] != 8; + if (err) + abort (); + } + return 0; +} Jakub