Hi! On Wed, 2 Dec 2015 16:58:45 +0100, I wrote: > Cesar and Jim copied, for help with Fortran and generally testsuite > things. > > On Mon, 30 Nov 2015 19:30:34 +0000, Julian Brown wrote: > > [patch] > > First, thanks! Aside from a number of formatting/re-ordering changes, the front end changes were basically still the same, but otherwise (middle end, libgomp) the patch as committed to trunk in r231118 was quite (totally?) ;-) different from the code we had on gomp-4_0-branch, so I had to spend some time on merging, cleaning things up. > What about the test cases present on gomp-4_0-branch, > gcc/testsuite/c-c++-common/goacc/host_data-1.c, > gcc/testsuite/c-c++-common/goacc/host_data-2.c, > gcc/testsuite/c-c++-common/goacc/host_data-3.c, and > gcc/testsuite/c-c++-common/goacc/host_data-4.c, [...] In the merge, I had to move two use_device usages from c-c++-common/goacc/host_data-1.c (was accepted) to c-c++-common/goacc/host_data-2.c (now rejected); I hope that's correct. > Your submission/commit didn't have any execution tests for OpenACC > host_data in Fortran. On gomp-4_0-branch, there is > libgomp/testsuite/libgomp.oacc-fortran/host_data-1.f90 at least. ..., but this one now FAILs (ICE) as follows: [...]/source-gcc/libgomp/testsuite/libgomp.oacc-fortran/host_data-1.f90:11:0: internal compiler error: in scan_omp_target, at omp-low.c:3218 0xa33e80 scan_omp_target [...]/source-gcc/gcc/omp-low.c:3218 0xa33e80 scan_omp_1_stmt [...]/source-gcc/gcc/omp-low.c:3980 0x8e4e7e walk_gimple_stmt(gimple_stmt_iterator*, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*) [...]/source-gcc/gcc/gimple-walk.c:555 0x8e50b8 walk_gimple_seq_mod(gimple**, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*) [...]/source-gcc/gcc/gimple-walk.c:51 0x8e4f62 walk_gimple_stmt(gimple_stmt_iterator*, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*) [...]/source-gcc/gcc/gimple-walk.c:583 0x8e50b8 walk_gimple_seq_mod(gimple**, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*) [...]/source-gcc/gcc/gimple-walk.c:51 0x8e4ff2 walk_gimple_stmt(gimple_stmt_iterator*, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*) [...]/source-gcc/gcc/gimple-walk.c:619 0x8e50b8 walk_gimple_seq_mod(gimple**, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*) [...]/source-gcc/gcc/gimple-walk.c:51 0xa02479 scan_omp [...]/source-gcc/gcc/omp-low.c:4024 0xa32ea5 scan_omp_target [...]/source-gcc/gcc/omp-low.c:3204 0xa32ea5 scan_omp_1_stmt [...]/source-gcc/gcc/omp-low.c:3980 0x8e4e7e walk_gimple_stmt(gimple_stmt_iterator*, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*) [...]/source-gcc/gcc/gimple-walk.c:555 0x8e50b8 walk_gimple_seq_mod(gimple**, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*) [...]/source-gcc/gcc/gimple-walk.c:51 0x8e4ff2 walk_gimple_stmt(gimple_stmt_iterator*, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*) [...]/source-gcc/gcc/gimple-walk.c:619 0x8e50b8 walk_gimple_seq_mod(gimple**, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*) [...]/source-gcc/gcc/gimple-walk.c:51 0x8e4f62 walk_gimple_stmt(gimple_stmt_iterator*, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*) [...]/source-gcc/gcc/gimple-walk.c:583 0x8e50b8 walk_gimple_seq_mod(gimple**, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*) [...]/source-gcc/gcc/gimple-walk.c:51 0xa02479 scan_omp [...]/source-gcc/gcc/omp-low.c:4024 0xa3f35a execute_lower_omp [...]/source-gcc/gcc/omp-low.c:16735 0xa3f35a execute [...]/source-gcc/gcc/omp-low.c:16782 Maybe that's due to the gcc/gimplify.c:gimplify_scan_omp_clauses issue mentioned in , or maybe something else? (XFAILed for now.) (For avoidance of doubt, the merge does not include my "Some OpenACC host_data cleanup" commit, trunk r231184, which will get merged into gomp-4_0-branch later.) So, merging trunk r231118 into gomp-4_0-branch, I effectively applied the following patch, in r231207. Please verify. For instance, do we need to re-instantiate any of the testsuite code that we've lost here, or is all of that actually not supported? commit 15723d76ae42dfe3f7201e0e3c6cbd9f4fc480b2 Merge: e08db3c 571b348 Author: tschwinge Date: Wed Dec 2 21:52:25 2015 +0000 svn merge -r 231117:231118 svn+ssh://gcc.gnu.org/svn/gcc/trunk git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@231207 138bc75d-0d04-0410-961f-82ee72b054a4 gcc/ChangeLog | 29 ++++ gcc/c-family/ChangeLog | 8 + gcc/c/ChangeLog | 14 ++ gcc/c/c-parser.c | 12 +- gcc/c/c-typeck.c | 2 +- gcc/cp/ChangeLog | 14 ++ gcc/cp/parser.c | 62 ++++---- gcc/cp/semantics.c | 6 +- gcc/gimple-pretty-print.c | 3 + gcc/gimple.h | 2 + gcc/gimplify.c | 177 +++++---------------- gcc/omp-builtins.def | 4 +- gcc/omp-low.c | 25 ++- gcc/testsuite/c-c++-common/goacc/host_data-1.c | 4 +- gcc/testsuite/c-c++-common/goacc/host_data-2.c | 10 ++ gcc/testsuite/gfortran.dg/goacc/coarray.f95 | 2 - gcc/tree-nested.c | 2 + libgomp/ChangeLog | 12 ++ libgomp/libgomp.map | 2 +- libgomp/oacc-mem.c | 32 ---- libgomp/oacc-parallel.c | 40 +++++ .../libgomp.oacc-c-c++-common/host_data-1.c | 39 +---- .../libgomp.oacc-c-c++-common/host_data-2.c | 57 +++---- .../libgomp.oacc-c-c++-common/host_data-3.c | 29 ++++ .../libgomp.oacc-c-c++-common/host_data-4.c | 29 ++++ .../libgomp.oacc-c-c++-common/host_data-5.c | 38 +++++ .../libgomp.oacc-c-c++-common/host_data-6.c | 31 ++++ .../testsuite/libgomp.oacc-fortran/host_data-1.f90 | 5 +- 28 files changed, 394 insertions(+), 296 deletions(-) [diff --git gcc/ChangeLog gcc/ChangeLog] [diff --git gcc/c-family/ChangeLog gcc/c-family/ChangeLog] [diff --git gcc/c/ChangeLog gcc/c/ChangeLog] diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 7191665..0251b80 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -10315,10 +10315,10 @@ c_parser_omp_clause_name (c_parser *parser, bool consume_token = true) result = PRAGMA_OMP_CLAUSE_UNIFORM; else if (!strcmp ("untied", p)) result = PRAGMA_OMP_CLAUSE_UNTIED; - else if (!strcmp ("use_device", p)) - result = PRAGMA_OACC_CLAUSE_USE_DEVICE; else if (!strcmp ("use_device_ptr", p)) result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR; + else if (!strcmp ("use_device", p)) + result = PRAGMA_OACC_CLAUSE_USE_DEVICE; break; case 'v': if (!strcmp ("vector", p)) @@ -13113,6 +13113,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_omp_clause_reduction (parser, clauses); c_name = "reduction"; break; + case PRAGMA_OACC_CLAUSE_USE_DEVICE: + clauses = c_parser_oacc_clause_use_device (parser, clauses); + c_name = "use_device"; + break; case PRAGMA_OACC_CLAUSE_SEQ: clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ, clauses); @@ -13122,10 +13126,6 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_oacc_clause_tile (parser, clauses); c_name = "tile"; break; - case PRAGMA_OACC_CLAUSE_USE_DEVICE: - clauses = c_parser_oacc_clause_use_device (parser, clauses); - c_name = "use_device"; - break; case PRAGMA_OACC_CLAUSE_VECTOR: c_name = "vector"; clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR, diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c index c40f6da..4659814 100644 --- gcc/c/c-typeck.c +++ gcc/c/c-typeck.c @@ -13168,6 +13168,7 @@ c_finish_omp_clauses (tree clauses, bool is_oacc, bool is_omp, bool declare_simd bitmap_set_bit (&map_head, DECL_UID (t)); goto check_dup_generic; + case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_USE_DEVICE_PTR: t = OMP_CLAUSE_DECL (c); @@ -13230,7 +13231,6 @@ c_finish_omp_clauses (tree clauses, bool is_oacc, bool is_omp, bool declare_simd case OMP_CLAUSE_GANG: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_BIND: case OMP_CLAUSE_NOHOST: case OMP_CLAUSE_TILE: [diff --git gcc/cp/ChangeLog gcc/cp/ChangeLog] diff --git gcc/cp/parser.c gcc/cp/parser.c index ac3f45c..d32aa91 100644 --- gcc/cp/parser.c +++ gcc/cp/parser.c @@ -29242,10 +29242,10 @@ cp_parser_omp_clause_name (cp_parser *parser, bool consume_token = true) result = PRAGMA_OMP_CLAUSE_UNIFORM; else if (!strcmp ("untied", p)) result = PRAGMA_OMP_CLAUSE_UNTIED; - else if (!strcmp ("use_device", p)) - result = PRAGMA_OACC_CLAUSE_USE_DEVICE; else if (!strcmp ("use_device_ptr", p)) result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR; + else if (!strcmp ("use_device", p)) + result = PRAGMA_OACC_CLAUSE_USE_DEVICE; break; case 'v': if (!strcmp ("vector", p)) @@ -31752,6 +31752,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_omp_clause_reduction (parser, clauses); c_name = "reduction"; break; + case PRAGMA_OACC_CLAUSE_USE_DEVICE: + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE, + clauses); + c_name = "use_device"; + break; case PRAGMA_OACC_CLAUSE_SEQ: clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ, clauses, here); @@ -31761,11 +31766,6 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_clause_tile (parser, here, clauses); c_name = "tile"; break; - case PRAGMA_OACC_CLAUSE_USE_DEVICE: - clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE, - clauses); - c_name = "use_device"; - break; case PRAGMA_OACC_CLAUSE_VECTOR: c_name = "vector"; clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR, @@ -34671,6 +34671,30 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok) return stmt; } +#define OACC_HOST_DATA_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) ) + +/* OpenACC 2.0: + # pragma acc host_data new-line + structured-block */ + +static tree +cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok) +{ + tree stmt, clauses, block; + unsigned int save; + + clauses = cp_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK, + "#pragma acc host_data", pragma_tok); + + block = begin_omp_parallel (); + save = cp_parser_begin_omp_structured_block (parser); + cp_parser_statement (parser, NULL_TREE, false, NULL); + cp_parser_end_omp_structured_block (parser, save); + stmt = finish_oacc_host_data (clauses, block); + return stmt; +} + /* OpenACC 2.0: # pragma acc declare oacc-data-clause[optseq] new-line */ @@ -34823,30 +34847,6 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) return NULL_TREE; } -#define OACC_HOST_DATA_CLAUSE_MASK \ - ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) ) - -/* OpenACC 2.0: - # pragma acc host_data new-line - structured-block */ - -static tree -cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok) -{ - tree stmt, clauses, block; - unsigned int save; - - clauses = cp_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK, - "#pragma acc host_data", pragma_tok); - - block = begin_omp_parallel (); - save = cp_parser_begin_omp_structured_block (parser); - cp_parser_statement (parser, NULL_TREE, false, NULL); - cp_parser_end_omp_structured_block (parser, save); - stmt = finish_oacc_host_data (clauses, block); - return stmt; -} - /* OpenACC 2.0: # pragma acc enter data oacc-enter-data-clause[optseq] new-line diff --git gcc/cp/semantics.c gcc/cp/semantics.c index 0d7e23d..a9a6671 100644 --- gcc/cp/semantics.c +++ gcc/cp/semantics.c @@ -6911,6 +6911,7 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields, } break; + case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_USE_DEVICE_PTR: field_ok = allow_fields; @@ -6948,7 +6949,6 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields, case OMP_CLAUSE_SIMD: case OMP_CLAUSE_DEFAULTMAP: case OMP_CLAUSE__CILK_FOR_COUNT_: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_AUTO: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_SEQ: @@ -7483,9 +7483,9 @@ tree finish_oacc_host_data (tree clauses, tree block) { tree stmt; - + block = finish_omp_structured_block (block); - + stmt = make_node (OACC_HOST_DATA); TREE_TYPE (stmt) = void_type_node; OACC_HOST_DATA_CLAUSES (stmt) = clauses; diff --git gcc/gimple-pretty-print.c gcc/gimple-pretty-print.c index 6c4e42c..c0f7c20 100644 --- gcc/gimple-pretty-print.c +++ gcc/gimple-pretty-print.c @@ -1356,6 +1356,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gomp_target *gs, case GF_OMP_TARGET_KIND_OACC_DECLARE: kind = " oacc_declare"; break; + case GF_OMP_TARGET_KIND_OACC_HOST_DATA: + kind = " oacc_host_data"; + break; default: gcc_unreachable (); } diff --git gcc/gimple.h gcc/gimple.h index 4c90bd7..7aaf785 100644 --- gcc/gimple.h +++ gcc/gimple.h @@ -171,6 +171,7 @@ enum gf_mask { GF_OMP_TARGET_KIND_OACC_UPDATE = 8, GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 9, GF_OMP_TARGET_KIND_OACC_DECLARE = 10, + GF_OMP_TARGET_KIND_OACC_HOST_DATA = 11, /* True on an GIMPLE_OMP_RETURN statement if the return does not require a thread synchronization via some sort of barrier. The exact barrier @@ -6006,6 +6007,7 @@ is_gimple_omp_oacc (const gimple *stmt) case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_DECLARE: + case GF_OMP_TARGET_KIND_OACC_HOST_DATA: return true; default: return false; diff --git gcc/gimplify.c gcc/gimplify.c index 3bb3bfe..b00de81 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -90,10 +90,8 @@ enum gimplify_omp_var_data /* Flag for shared vars that are or might be stored to in the region. */ GOVD_WRITTEN = 131072, - GOVD_USE_DEVICE = 1 << 18, - /* OpenACC deviceptr clause. */ - GOVD_USE_DEVPTR = 1 << 19, + GOVD_USE_DEVPTR = 1 << 18, GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR @@ -122,18 +120,16 @@ enum omp_region_type ORT_TARGET = 0x20, ORT_COMBINED_TARGET = 0x21, - ORT_HOST_DATA = 0x40, - /* OpenACC variants. */ - ORT_ACC = 0x80, /* A generic OpenACC region. */ + ORT_ACC = 0x40, /* A generic OpenACC region. */ ORT_ACC_DATA = ORT_ACC | ORT_TARGET_DATA, /* Data construct. */ ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET, /* Parallel construct */ - ORT_ACC_KERNELS = ORT_ACC | ORT_TARGET | 0x100, /* Kernels construct. */ - ORT_ACC_HOST = ORT_ACC | ORT_HOST_DATA, + ORT_ACC_KERNELS = ORT_ACC | ORT_TARGET | 0x80, /* Kernels construct. */ + ORT_ACC_HOST_DATA = ORT_ACC | ORT_TARGET_DATA | 0x80, /* Host data. */ /* Dummy OpenMP region, used to disable expansion of DECL_VALUE_EXPRs in taskloop pre body. */ - ORT_NONE = 0x200 + ORT_NONE = 0x100 }; /* Gimplify hashtable helper. */ @@ -6126,8 +6122,6 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) for (; octx; octx = octx->outer_context) { - if (octx->region_type & ORT_HOST_DATA) - continue; if (!(octx->region_type & (ORT_TARGET_DATA | ORT_TARGET))) break; splay_tree_node n2 @@ -6135,6 +6129,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) (splay_tree_key) decl); if (n2) { + if (octx->region_type == ORT_ACC_HOST_DATA) + error ("variable %qE declared in enclosing " + "% region", DECL_NAME (decl)); nflags |= GOVD_MAP; goto found_outer; } @@ -6436,6 +6433,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_TARGET_DATA: case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: + case OACC_HOST_DATA: ctx->target_firstprivatize_array_bases = true; default: break; @@ -6571,10 +6569,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, || outer_ctx->region_type == ORT_ACC_DATA)) redvec.safe_push (OMP_CLAUSE_DECL (c)); goto do_add_decl; - case OMP_CLAUSE_USE_DEVICE: - flags = GOVD_USE_DEVICE | GOVD_EXPLICIT; - check_non_private = "use_device"; - goto do_add; case OMP_CLAUSE_LINEAR: if (gimplify_expr (&OMP_CLAUSE_LINEAR_STEP (c), pre_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR) @@ -6709,6 +6703,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_TARGET_DATA: case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: + case OACC_HOST_DATA: if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) @@ -6721,6 +6716,22 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } if (remove) break; + if (DECL_P (decl) && outer_ctx && (region_type & ORT_ACC)) + { + struct gimplify_omp_ctx *octx; + for (octx = outer_ctx; octx; octx = octx->outer_context) + { + if (octx->region_type != ORT_ACC_HOST_DATA) + break; + splay_tree_node n2 + = splay_tree_lookup (octx->variables, + (splay_tree_key) decl); + if (n2) + error_at (OMP_CLAUSE_LOCATION (c), "variable %qE " + "declared in enclosing % region", + DECL_NAME (decl)); + } + } if (OMP_CLAUSE_SIZE (c) == NULL_TREE) OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl) : TYPE_SIZE_UNIT (TREE_TYPE (decl)); @@ -7120,6 +7131,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } goto do_notice; + case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT; goto do_add; @@ -7639,7 +7651,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) code = OMP_CLAUSE_FIRSTPRIVATE; else if (flags & GOVD_LASTPRIVATE) code = OMP_CLAUSE_LASTPRIVATE; - else if (flags & (GOVD_ALIGNED | GOVD_USE_DEVICE)) + else if (flags & GOVD_ALIGNED) return 0; else gcc_unreachable (); @@ -8244,126 +8256,6 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p) *expr_p = NULL_TREE; } -static tree -gimplify_oacc_host_data_1 (tree *tp, int *walk_subtrees, void *data ATTRIBUTE_UNUSED) -{ - splay_tree_node n = NULL; - location_t loc = EXPR_LOCATION (*tp); - - switch (TREE_CODE (*tp)) - { - case ADDR_EXPR: - { - tree decl = TREE_OPERAND (*tp, 0); - - switch (TREE_CODE (decl)) - { - case ARRAY_REF: - case ARRAY_RANGE_REF: - case COMPONENT_REF: - case VIEW_CONVERT_EXPR: - case REALPART_EXPR: - case IMAGPART_EXPR: - if (TREE_CODE (TREE_OPERAND (decl, 0)) == VAR_DECL) - n = splay_tree_lookup (gimplify_omp_ctxp->variables, - (splay_tree_key) TREE_OPERAND (decl, 0)); - break; - - case VAR_DECL: - n = splay_tree_lookup (gimplify_omp_ctxp->variables, - (splay_tree_key) decl); - break; - - default: - ; - } - - if (n != NULL && (n->value & GOVD_USE_DEVICE) != 0) - { - tree t = builtin_decl_explicit (BUILT_IN_GOACC_DEVICEPTR); - *tp = build_call_expr_loc (loc, t, 1, *tp); - } - - *walk_subtrees = 0; - } - break; - - case VAR_DECL: - { - tree decl = *tp; - - n = splay_tree_lookup (gimplify_omp_ctxp->variables, - (splay_tree_key) decl); - - if (n != NULL && (n->value & GOVD_USE_DEVICE) != 0) - { - if (!POINTER_TYPE_P (TREE_TYPE (decl))) - return decl; - - tree t = builtin_decl_explicit (BUILT_IN_GOACC_DEVICEPTR); - *tp = build_call_expr_loc (loc, t, 1, *tp); - *walk_subtrees = 0; - } - } - break; - - case OACC_PARALLEL: - case OACC_KERNELS: - case OACC_LOOP: - *walk_subtrees = 0; - break; - - default: - ; - } - - return NULL_TREE; -} - -static enum gimplify_status -gimplify_oacc_host_data (tree *expr_p, gimple_seq *pre_p) -{ - tree expr = *expr_p, orig_body; - gimple_seq body = NULL; - - gimplify_scan_omp_clauses (&OACC_HOST_DATA_CLAUSES (expr), pre_p, - ORT_ACC_HOST, OACC_HOST_DATA); - - orig_body = OACC_HOST_DATA_BODY (expr); - - /* Perform a pre-pass over the host_data region's body, inserting calls to - GOACC_deviceptr where appropriate. */ - - tree ret = walk_tree_without_duplicates (&orig_body, - &gimplify_oacc_host_data_1, 0); - - if (ret) - { - error_at (EXPR_LOCATION (expr), - "undefined use of variable %qE in host_data region", - DECL_NAME (ret)); - gimplify_adjust_omp_clauses (pre_p, body, &OACC_HOST_DATA_CLAUSES (expr), - OACC_HOST_DATA); - return GS_ERROR; - } - - push_gimplify_context (); - - gimple *g = gimplify_and_return_first (orig_body, &body); - - if (gimple_code (g) == GIMPLE_BIND) - pop_gimplify_context (g); - else - pop_gimplify_context (NULL); - - gimplify_adjust_omp_clauses (pre_p, body, &OACC_HOST_DATA_CLAUSES (expr), - OACC_HOST_DATA); - - gimplify_seq_add_stmt (pre_p, g); - - return GS_ALL_DONE; -} - /* Gimplify the contents of an OMP_PARALLEL statement. This involves gimplification of the body, as well as scanning the body for used variables. We need to do this scan now, because variable-sized @@ -9648,6 +9540,9 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) case OMP_TEAMS: ort = OMP_TEAMS_COMBINED (expr) ? ORT_COMBINED_TEAMS : ORT_TEAMS; break; + case OACC_HOST_DATA: + ort = ORT_ACC_HOST_DATA; + break; default: gcc_unreachable (); } @@ -9673,6 +9568,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) switch (TREE_CODE (expr)) { case OACC_DATA: + case OACC_HOST_DATA: end_ix = BUILT_IN_GOACC_DATA_END; break; case OMP_TARGET_DATA: @@ -9705,6 +9601,10 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_KERNELS, OMP_CLAUSES (expr)); break; + case OACC_HOST_DATA: + stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_HOST_DATA, + OMP_CLAUSES (expr)); + break; case OACC_PARALLEL: stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_PARALLEL, OMP_CLAUSES (expr)); @@ -10814,15 +10714,12 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, ret = GS_ALL_DONE; break; - case OACC_HOST_DATA: - ret = gimplify_oacc_host_data (expr_p, pre_p); - break; - case OACC_DECLARE: gimplify_oacc_declare (expr_p, pre_p); ret = GS_ALL_DONE; break; + case OACC_HOST_DATA: case OACC_DATA: case OACC_KERNELS: case OACC_PARALLEL: diff --git gcc/omp-builtins.def gcc/omp-builtins.def index 63e5e6e..35f5014 100644 --- gcc/omp-builtins.def +++ gcc/omp-builtins.def @@ -47,8 +47,8 @@ DEF_GOACC_BUILTIN (BUILT_IN_GOACC_UPDATE, "GOACC_update", DEF_GOACC_BUILTIN (BUILT_IN_GOACC_WAIT, "GOACC_wait", BT_FN_VOID_INT_INT_VAR, ATTR_NOTHROW_LIST) -DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DEVICEPTR, "GOACC_deviceptr", - BT_FN_PTR_PTR, ATTR_CONST_NOTHROW_LEAF_LIST) +DEF_GOACC_BUILTIN (BUILT_IN_GOACC_HOST_DATA, "GOACC_host_data", + BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST) DEF_GOACC_BUILTIN_COMPILER (BUILT_IN_ACC_ON_DEVICE, "acc_on_device", BT_FN_INT_INT, ATTR_CONST_NOTHROW_LEAF_LIST) diff --git gcc/omp-low.c gcc/omp-low.c index a1e7a14..88e41b8 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -2071,6 +2071,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) } break; + case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: decl = OMP_CLAUSE_DECL (c); if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) @@ -2274,7 +2275,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) break; case OMP_CLAUSE_DEVICE_RESIDENT: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE__CACHE_: sorry ("Clause not supported yet"); break; @@ -2430,6 +2430,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_SIMD: case OMP_CLAUSE_NOGROUP: case OMP_CLAUSE_DEFAULTMAP: + case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE__CILK_FOR_COUNT_: case OMP_CLAUSE_ASYNC: @@ -2448,7 +2449,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) break; case OMP_CLAUSE_DEVICE_RESIDENT: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE__CACHE_: sorry ("Clause not supported yet"); break; @@ -3763,6 +3763,8 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) case GF_OMP_TARGET_KIND_OACC_UPDATE: stmt_name = "update"; break; case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: stmt_name = "enter/exit data"; break; + case GF_OMP_TARGET_KIND_OACC_HOST_DATA: stmt_name = "host_data"; + break; default: gcc_unreachable (); } switch (gimple_omp_target_kind (ctx->stmt)) @@ -3774,6 +3776,8 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) case GF_OMP_TARGET_KIND_OACC_KERNELS: ctx_stmt_name = "kernels"; break; case GF_OMP_TARGET_KIND_OACC_DATA: ctx_stmt_name = "data"; break; + case GF_OMP_TARGET_KIND_OACC_HOST_DATA: + ctx_stmt_name = "host_data"; break; default: gcc_unreachable (); } @@ -12730,6 +12734,7 @@ expand_omp_target (struct omp_region *region) break; case GF_OMP_TARGET_KIND_DATA: case GF_OMP_TARGET_KIND_OACC_DATA: + case GF_OMP_TARGET_KIND_OACC_HOST_DATA: data_region = true; break; default: @@ -12980,6 +12985,9 @@ expand_omp_target (struct omp_region *region) case GF_OMP_TARGET_KIND_OACC_DECLARE: start_ix = BUILT_IN_GOACC_DECLARE; break; + case GF_OMP_TARGET_KIND_OACC_HOST_DATA: + start_ix = BUILT_IN_GOACC_HOST_DATA; + break; default: gcc_unreachable (); } @@ -13104,6 +13112,7 @@ expand_omp_target (struct omp_region *region) case BUILT_IN_GOACC_DATA_START: case BUILT_IN_GOACC_DECLARE: case BUILT_IN_GOMP_TARGET_DATA: + case BUILT_IN_GOACC_HOST_DATA: break; case BUILT_IN_GOMP_TARGET: case BUILT_IN_GOMP_TARGET_UPDATE: @@ -13445,6 +13454,7 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent, case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: case GF_OMP_TARGET_KIND_OACC_DATA: + case GF_OMP_TARGET_KIND_OACC_HOST_DATA: if (is_gimple_omp_oacc (stmt)) region->kind = gimple_omp_target_kind (stmt); break; @@ -15277,6 +15287,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) break; case GF_OMP_TARGET_KIND_DATA: case GF_OMP_TARGET_KIND_OACC_DATA: + case GF_OMP_TARGET_KIND_OACC_HOST_DATA: data_region = true; break; default: @@ -15485,6 +15496,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) } break; + case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: var = OMP_CLAUSE_DECL (c); @@ -15870,12 +15882,14 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) build_int_cstu (tkind_type, tkind)); break; + case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: ovar = OMP_CLAUSE_DECL (c); var = lookup_decl_in_outer_ctx (ovar, ctx); x = build_sender_ref (ovar, ctx); - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE) tkind = GOMP_MAP_USE_DEVICE_PTR; else tkind = GOMP_MAP_FIRSTPRIVATE_INT; @@ -16078,10 +16092,12 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_build_assign (new_var, x)); } break; + case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: var = OMP_CLAUSE_DECL (c); - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE) x = build_sender_ref (var, ctx); else x = build_receiver_ref (var, false, ctx); @@ -17076,6 +17092,7 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region, case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: case GF_OMP_TARGET_KIND_OACC_DATA: + case GF_OMP_TARGET_KIND_OACC_HOST_DATA: break; case GF_OMP_TARGET_KIND_UPDATE: case GF_OMP_TARGET_KIND_ENTER_DATA: diff --git gcc/testsuite/c-c++-common/goacc/host_data-1.c gcc/testsuite/c-c++-common/goacc/host_data-1.c index 521c854..a8922df 100644 --- gcc/testsuite/c-c++-common/goacc/host_data-1.c +++ gcc/testsuite/c-c++-common/goacc/host_data-1.c @@ -1,13 +1,11 @@ /* Test valid use of host_data directive. */ /* { dg-do compile } */ -int v0; int v1[3][3]; void f (void) { - int v2 = 3; -#pragma acc host_data use_device(v2, v0, v1) +#pragma acc host_data use_device(v1) ; } diff --git gcc/testsuite/c-c++-common/goacc/host_data-2.c gcc/testsuite/c-c++-common/goacc/host_data-2.c index e5213a0..1dd5be7 100644 --- gcc/testsuite/c-c++-common/goacc/host_data-2.c +++ gcc/testsuite/c-c++-common/goacc/host_data-2.c @@ -10,4 +10,14 @@ f (void) int v2 = 3; #pragma acc host_data copy(v2) /* { dg-error "not valid for" } */ ; + +#pragma acc host_data use_device(v2) + ; + /* { dg-error ".use_device. variable is neither a pointer nor an array" "" { target c } 14 } */ + /* { dg-error ".use_device. variable is neither a pointer, nor an arraynor reference to pointer or array" "" { target c++ } 14 } */ + +#pragma acc host_data use_device(v0) + ; + /* { dg-error ".use_device. variable is neither a pointer nor an array" "" { target c } 19 } */ + /* { dg-error ".use_device. variable is neither a pointer, nor an arraynor reference to pointer or array" "" { target c++ } 19 } */ } diff --git gcc/testsuite/gfortran.dg/goacc/coarray.f95 gcc/testsuite/gfortran.dg/goacc/coarray.f95 index 0ca14e2..d2f10d5 100644 --- gcc/testsuite/gfortran.dg/goacc/coarray.f95 +++ gcc/testsuite/gfortran.dg/goacc/coarray.f95 @@ -5,8 +5,6 @@ ! { dg-xfail-if "" { *-*-* } } ! { dg-excess-errors "TODO" } -! TODO: These cases must fail - module test contains subroutine oacc1(a) diff --git gcc/tree-nested.c gcc/tree-nested.c index 8b5aba2..da19e8d 100644 --- gcc/tree-nested.c +++ gcc/tree-nested.c @@ -1072,6 +1072,7 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_SHARED: case OMP_CLAUSE_TO_DECLARE: case OMP_CLAUSE_LINK: + case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: do_decl_clause: @@ -1743,6 +1744,7 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_SHARED: case OMP_CLAUSE_TO_DECLARE: case OMP_CLAUSE_LINK: + case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: do_decl_clause: [diff --git libgomp/ChangeLog libgomp/ChangeLog] diff --git libgomp/libgomp.map libgomp/libgomp.map index cceb92d..a42142f 100644 --- libgomp/libgomp.map +++ libgomp/libgomp.map @@ -394,11 +394,11 @@ GOACC_2.0.1 { global: GOACC_declare; GOACC_parallel_keyed; + GOACC_host_data; } GOACC_2.0; GOACC_2.0.GOMP_4_BRANCH { global: - GOACC_deviceptr; GOMP_set_offload_targets; } GOACC_2.0.1; diff --git libgomp/oacc-mem.c libgomp/oacc-mem.c index 11edcce..588782b 100644 --- libgomp/oacc-mem.c +++ libgomp/oacc-mem.c @@ -203,38 +203,6 @@ acc_deviceptr (void *h) return d; } -/* This function is used as a helper in generated code to implement pointer - lookup in host_data regions. Unlike acc_deviceptr, it returns its argument - unchanged on a shared-memory system (e.g. the host). */ - -void * -GOACC_deviceptr (void *h) -{ - splay_tree_key n; - void *d; - void *offset; - - goacc_lazy_initialize (); - - struct goacc_thread *thr = goacc_thread (); - - if ((thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) == 0) - { - n = lookup_host (thr->dev, h, 1); - - if (!n) - return NULL; - - offset = h - n->host_start; - - d = n->tgt->tgt_start + n->tgt_offset + offset; - - return d; - } - else - return h; -} - /* Return the host pointer that corresponds to device data D. Or NULL if no mapping. */ diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c index d66e343..e60a61b 100644 --- libgomp/oacc-parallel.c +++ libgomp/oacc-parallel.c @@ -555,6 +555,46 @@ GOACC_wait (int async, int num_waits, ...) goacc_thread ()->dev->openacc.async_wait_all_async_func (acc_async_noval); } +void +GOACC_host_data (int device, size_t mapnum, + void **hostaddrs, size_t *sizes, unsigned short *kinds) +{ + bool host_fallback = device == GOMP_DEVICE_HOST_FALLBACK; + struct target_mem_desc *tgt; + +#ifdef HAVE_INTTYPES_H + gomp_debug (0, "%s: mapnum=%"PRIu64", hostaddrs=%p, size=%p, kinds=%p\n", + __FUNCTION__, (uint64_t) mapnum, hostaddrs, sizes, kinds); +#else + gomp_debug (0, "%s: mapnum=%lu, hostaddrs=%p, sizes=%p, kinds=%p\n", + __FUNCTION__, (unsigned long) mapnum, hostaddrs, sizes, kinds); +#endif + + goacc_lazy_initialize (); + + struct goacc_thread *thr = goacc_thread (); + struct gomp_device_descr *acc_dev = thr->dev; + + /* Host fallback or 'do nothing'. */ + if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + || host_fallback) + { + tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, + GOMP_MAP_VARS_OPENACC); + tgt->prev = thr->mapped_data; + thr->mapped_data = tgt; + + return; + } + + gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__); + tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true, + GOMP_MAP_VARS_OPENACC); + gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__); + tgt->prev = thr->mapped_data; + thr->mapped_data = tgt; +} + int GOACC_get_num_threads (int gang, int worker, int vector) { diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c index 15ccb27..51745ba 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c @@ -1,7 +1,6 @@ /* { dg-do run { target openacc_nvidia_accel_selected } } */ /* { dg-additional-options "-lcuda -lcublas -lcudart" } */ -#include #include #include #include @@ -30,35 +29,13 @@ saxpy_target (int n, float a, float *x, float *y) int main(int argc, char **argv) { - const int N = 8; +#define N 8 int i; - float *x_ref, *y_ref; - float *x, *y; + float x_ref[N], y_ref[N]; + float x[N], y[N]; cublasHandle_t h; float a = 2.0; - x_ref = (float*) malloc (N * sizeof(float)); - y_ref = (float*) malloc (N * sizeof(float)); - - x = (float*) malloc (N * sizeof(float)); - y = (float*) malloc (N * sizeof(float)); - -#pragma acc data copyin (x[0:N]) copy (y[0:N]) - { - float *xp, *yp; -#pragma acc host_data use_device (x, y) - { -#pragma acc parallel pcopy (xp, yp) present (x, y) - { - xp = x; - yp = y; - } - } - - if (xp != acc_deviceptr (x) || yp != acc_deviceptr (y)) - abort (); - } - for (i = 0; i < N; i++) { x[i] = x_ref[i] = 4.0 + i; @@ -106,13 +83,11 @@ main(int argc, char **argv) for (i = 0; i < N; i++) y[i] = 3.0; -#pragma acc data copyin (x[0:N]) copyin (a, N) copy (y[0:N]) + /* There's no need to use host_data here. */ +#pragma acc data copyin (x[0:N]) copyin (a) copy (y[0:N]) { -#pragma acc host_data use_device (x, y) - { -#pragma acc parallel present (x[0:N]) pcopy (y[0:N]) present (a, N) - saxpy_target (N, a, x, y); - } +#pragma acc parallel present (x[0:N]) pcopy (y[0:N]) present (a) + saxpy_target (N, a, x, y); } for (i = 0; i < N; i++) diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c index 511ec64..614f143 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c @@ -1,50 +1,31 @@ -/* { dg-do run } */ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include +#include -struct by_lightning { - int a; - int b; - int c; -}; +char *global_in_host; -int main (int argc, char* argv[]) +void foo (char *in) { - int x; - void *q = NULL, *r = NULL, *p = NULL, *s = NULL, *t = NULL; - long u; - struct by_lightning on_the_head = {1, 2, 3}; - int arr[10], *f = NULL; - _Complex float cf; - #pragma acc enter data copyin (x, arr, on_the_head, cf) - #pragma acc host_data use_device (x, arr, on_the_head, cf) + if (!acc_is_present (global_in_host, sizeof (*global_in_host)) + || in != acc_deviceptr (global_in_host)) + abort (); +} + +int +main (int argc, char **argv) +{ + char mydata[1024]; + + global_in_host = mydata; + +#pragma acc data copyin(mydata) { - q = &x; +#pragma acc host_data use_device (mydata) { - f = &arr[5]; - r = f; - s = &__real__ cf; - t = &on_the_head.c; - u = (long) &__imag__ cf; - #pragma acc parallel copyout(p) present (x, arr, on_the_head, cf) - { - /* This will not (and must not) call GOACC_deviceptr, but '&x' will be - the address on the device (if appropriate) regardless. */ - p = &x; - } + foo (mydata); } } - #pragma acc exit data delete (x) - -#if ACC_MEM_SHARED - if (q != &x || f != &arr[5] || r != f || s != &(__real__ cf) - || t != &on_the_head.c || u != (long) &(__imag__ cf) || p != &x) - abort (); -#else - if (q == &x || f == &arr[5] || r != f || s == &(__real__ cf) - || t == &on_the_head.c || u == (long) &(__imag__ cf) || p == &x) - abort (); -#endif return 0; } diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-3.c new file mode 100644 index 0000000..7d9b5f7 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-3.c @@ -0,0 +1,29 @@ +/* { dg-do compile } */ + +#include +#include + +#define N 1024 + +int main (int argc, char* argv[]) +{ + int x[N]; + +#pragma acc data copyin (x[0:N]) + { + int *xp; +#pragma acc host_data use_device (x) + { + /* This use of the present clause is undefined behaviour for OpenACC. */ +#pragma acc parallel present (x) copyout (xp) /* { dg-error "variable 'x' declared in enclosing 'host_data' region" } */ + { + xp = x; + } + } + + if (xp != acc_deviceptr (x)) + abort (); + } + + return 0; +} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-4.c libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-4.c new file mode 100644 index 0000000..0ab5a35 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-4.c @@ -0,0 +1,29 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +#include +#include + +#define N 1024 + +int main (int argc, char* argv[]) +{ + int x[N], *xp2; + +#pragma acc data copyin (x[0:N]) + { + int *xp; +#pragma acc host_data use_device (x) + { +#pragma acc data + { + xp = x; + } + xp2 = x; + } + + if (xp != acc_deviceptr (x) || xp2 != xp) + abort (); + } + + return 0; +} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-5.c libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-5.c new file mode 100644 index 0000000..a3737a7 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-5.c @@ -0,0 +1,38 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +#include +#include + +#define N 1024 + +int main (int argc, char* argv[]) +{ + int x[N], y[N], *yp; + + yp = y + 1; + +#pragma acc data copyin (x[0:N]) + { + int *xp, *yp2; +#pragma acc host_data use_device (x) + { +#pragma acc data copyin (y) + { +#pragma acc host_data use_device (yp) + { + xp = x; + yp2 = yp; + } + + if (yp2 != acc_deviceptr (yp)) + abort (); + } + } + + if (xp != acc_deviceptr (x)) + abort (); + + } + + return 0; +} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c new file mode 100644 index 0000000..a841488 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c @@ -0,0 +1,31 @@ +/* { dg-do compile } */ + +#include +#include + +#define N 1024 + +int main (int argc, char* argv[]) +{ + int x[N]; + +#pragma acc data copyin (x[0:N]) + { + int *xp; +#pragma acc host_data use_device (x) + { + /* Here 'x' being implicitly firstprivate for the parallel region + conflicts with it being declared as use_device in the enclosing + host_data region. */ +#pragma acc parallel copyout (xp) + { + xp = x; /* { dg-error "variable 'x' declared in enclosing 'host_data' region" } */ + } + } + + if (xp != acc_deviceptr (x)) + abort (); + } + + return 0; +} diff --git libgomp/testsuite/libgomp.oacc-fortran/host_data-1.f90 libgomp/testsuite/libgomp.oacc-fortran/host_data-1.f90 index a219eaf..9bb79c3 100644 --- libgomp/testsuite/libgomp.oacc-fortran/host_data-1.f90 +++ libgomp/testsuite/libgomp.oacc-fortran/host_data-1.f90 @@ -1,6 +1,9 @@ -! { dg-do run } */ +! { dg-do run } ! { dg-additional-options "-cpp" } +! { dg-xfail-if "TODO" { *-*-* } } +! { dg-excess-errors "TODO" } + program test implicit none Grüße Thomas