diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 69fe3a7273e..02c2117836d 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,23 @@ +2015-12-15 Ilya Verbin + + * cgraphunit.c (output_in_order): Do not assemble "omp declare target + link" variables in ACCEL_COMPILER. + * gimplify.c (gimplify_adjust_omp_clauses): Do not remove mapping of + "omp declare target link" variables. + * omp-low.c (scan_sharing_clauses): Do not remove mapping of "omp + declare target link" variables. + (add_decls_addresses_to_decl_constructor): For "omp declare target link" + variables output address of the artificial pointer instead of address of + the variable. Set most significant bit of the size to mark them. + (pass_data_omp_target_link): New pass_data. + (pass_omp_target_link): New class. + (find_link_var_op): New static function. + (make_pass_omp_target_link): New function. + * passes.def: Add pass_omp_target_link. + * tree-pass.h (make_pass_omp_target_link): Declare. + * varpool.c (symbol_table::output_variables): Do not assemble "omp + declare target link" variables in ACCEL_COMPILER. + 2015-12-15 Bernd Schmidt PR middle-end/21273 diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog index db9a279c20f..525cc16f17b 100644 --- a/gcc/c-family/ChangeLog +++ b/gcc/c-family/ChangeLog @@ -1,3 +1,8 @@ +2015-12-15 Ilya Verbin + + * c-common.c (c_common_attribute_table): Handle "omp declare target + link" attribute. + 2015-12-14 Jakub Jelinek PR c/68833 diff --git a/gcc/c-family/c-common.c b/gcc/c-family/c-common.c index 9bc02fcf85c..4250cdf1cee 100644 --- a/gcc/c-family/c-common.c +++ b/gcc/c-family/c-common.c @@ -821,6 +821,8 @@ const struct attribute_spec c_common_attribute_table[] = handle_simd_attribute, false }, { "omp declare target", 0, 0, true, false, false, handle_omp_declare_target_attribute, false }, + { "omp declare target link", 0, 0, true, false, false, + handle_omp_declare_target_attribute, false }, { "alloc_align", 1, 1, false, true, true, handle_alloc_align_attribute, false }, { "assume_aligned", 1, 2, false, true, true, diff --git a/gcc/cgraphunit.c b/gcc/cgraphunit.c index 3d86c36488e..8443cb092a7 100644 --- a/gcc/cgraphunit.c +++ b/gcc/cgraphunit.c @@ -2210,6 +2210,13 @@ output_in_order (bool no_reorder) break; case ORDER_VAR: +#ifdef ACCEL_COMPILER + /* Do not assemble "omp declare target link" vars. */ + if (DECL_HAS_VALUE_EXPR_P (nodes[i].u.v->decl) + && lookup_attribute ("omp declare target link", + DECL_ATTRIBUTES (nodes[i].u.v->decl))) + break; +#endif nodes[i].u.v->assemble_decl (); break; diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 80c6bf2b90f..438efba57e1 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -7910,7 +7910,9 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); if ((ctx->region_type & ORT_TARGET) != 0 && !(n->value & GOVD_SEEN) - && GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) == 0) + && GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) == 0 + && !lookup_attribute ("omp declare target link", + DECL_ATTRIBUTES (decl))) { remove = true; /* For struct element mapping, if struct is never referenced diff --git a/gcc/lto/ChangeLog b/gcc/lto/ChangeLog index 6e905276d31..ac20a3f2dde 100644 --- a/gcc/lto/ChangeLog +++ b/gcc/lto/ChangeLog @@ -1,3 +1,9 @@ +2015-12-15 Ilya Verbin + + * lto.c: Include stringpool.h and fold-const.h. + (offload_handle_link_vars): New static function. + (lto_main): Call offload_handle_link_vars. + 2015-12-10 Jan Hubicka * lto.c (lto_read_in_decl_state): Unpickle compressed bit. diff --git a/gcc/lto/lto.c b/gcc/lto/lto.c index fcf7caf460a..5fd50dc1d26 100644 --- a/gcc/lto/lto.c +++ b/gcc/lto/lto.c @@ -50,6 +50,8 @@ along with GCC; see the file COPYING3. If not see #include "ipa-utils.h" #include "gomp-constants.h" #include "lto-symtab.h" +#include "stringpool.h" +#include "fold-const.h" /* Number of parallel tasks to run, -1 if we want to use GNU Make jobserver. */ @@ -3226,6 +3228,37 @@ lto_init (void) #endif } +/* Create artificial pointers for "omp declare target link" vars. */ + +static void +offload_handle_link_vars (void) +{ +#ifdef ACCEL_COMPILER + varpool_node *var; + FOR_EACH_VARIABLE (var) + if (lookup_attribute ("omp declare target link", + DECL_ATTRIBUTES (var->decl))) + { + tree type = build_pointer_type (TREE_TYPE (var->decl)); + tree link_ptr_var = make_node (VAR_DECL); + TREE_TYPE (link_ptr_var) = type; + TREE_USED (link_ptr_var) = 1; + TREE_STATIC (link_ptr_var) = 1; + DECL_MODE (link_ptr_var) = TYPE_MODE (type); + DECL_SIZE (link_ptr_var) = TYPE_SIZE (type); + DECL_SIZE_UNIT (link_ptr_var) = TYPE_SIZE_UNIT (type); + DECL_ARTIFICIAL (link_ptr_var) = 1; + tree var_name = DECL_ASSEMBLER_NAME (var->decl); + char *new_name + = ACONCAT ((IDENTIFIER_POINTER (var_name), "_linkptr", NULL)); + DECL_NAME (link_ptr_var) = get_identifier (new_name); + SET_DECL_ASSEMBLER_NAME (link_ptr_var, DECL_NAME (link_ptr_var)); + SET_DECL_VALUE_EXPR (var->decl, build_simple_mem_ref (link_ptr_var)); + DECL_HAS_VALUE_EXPR_P (var->decl) = 1; + } +#endif +} + /* Main entry point for the GIMPLE front end. This front end has three main personalities: @@ -3274,6 +3307,8 @@ lto_main (void) if (!seen_error ()) { + offload_handle_link_vars (); + /* If WPA is enabled analyze the whole call graph and create an optimization plan. Otherwise, read in all the function bodies and continue with optimization. */ diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 56434805cfe..676b1dfce88 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -2026,7 +2026,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, decl = OMP_CLAUSE_DECL (c); /* Global variables with "omp declare target" attribute don't need to be copied, the receiver side will use them - directly. */ + directly. However, global variables with "omp declare target link" + attribute need to be copied. */ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && DECL_P (decl) && ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER @@ -2034,7 +2035,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, != GOMP_MAP_FIRSTPRIVATE_REFERENCE)) || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) - && varpool_node::get_create (decl)->offloadable) + && varpool_node::get_create (decl)->offloadable + && !lookup_attribute ("omp declare target link", + DECL_ATTRIBUTES (decl))) break; if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER) @@ -18588,13 +18591,45 @@ add_decls_addresses_to_decl_constructor (vec *v_decls, for (unsigned i = 0; i < len; i++) { tree it = (*v_decls)[i]; - bool is_function = TREE_CODE (it) != VAR_DECL; + bool is_var = TREE_CODE (it) == VAR_DECL; + bool is_link_var + = is_var +#ifdef ACCEL_COMPILER + && DECL_HAS_VALUE_EXPR_P (it) +#endif + && lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (it)); - CONSTRUCTOR_APPEND_ELT (v_ctor, NULL_TREE, build_fold_addr_expr (it)); - if (!is_function) - CONSTRUCTOR_APPEND_ELT (v_ctor, NULL_TREE, - fold_convert (const_ptr_type_node, - DECL_SIZE_UNIT (it))); + tree size = NULL_TREE; + if (is_var) + size = fold_convert (const_ptr_type_node, DECL_SIZE_UNIT (it)); + + tree addr; + if (!is_link_var) + addr = build_fold_addr_expr (it); + else + { +#ifdef ACCEL_COMPILER + /* For "omp declare target link" vars add address of the pointer to + the target table, instead of address of the var. */ + tree value_expr = DECL_VALUE_EXPR (it); + tree link_ptr_decl = TREE_OPERAND (value_expr, 0); + varpool_node::finalize_decl (link_ptr_decl); + addr = build_fold_addr_expr (link_ptr_decl); +#else + addr = build_fold_addr_expr (it); +#endif + + /* Most significant bit of the size marks "omp declare target link" + vars in host and target tables. */ + unsigned HOST_WIDE_INT isize = tree_to_uhwi (size); + isize |= 1ULL << (int_size_in_bytes (const_ptr_type_node) + * BITS_PER_UNIT - 1); + size = wide_int_to_tree (const_ptr_type_node, isize); + } + + CONSTRUCTOR_APPEND_ELT (v_ctor, NULL_TREE, addr); + if (is_var) + CONSTRUCTOR_APPEND_ELT (v_ctor, NULL_TREE, size); } } @@ -19831,4 +19866,84 @@ make_pass_oacc_device_lower (gcc::context *ctxt) return new pass_oacc_device_lower (ctxt); } +/* "omp declare target link" handling pass. */ + +namespace { + +const pass_data pass_data_omp_target_link = +{ + GIMPLE_PASS, /* type */ + "omptargetlink", /* name */ + OPTGROUP_NONE, /* optinfo_flags */ + TV_NONE, /* tv_id */ + PROP_ssa, /* properties_required */ + 0, /* properties_provided */ + 0, /* properties_destroyed */ + 0, /* todo_flags_start */ + TODO_update_ssa, /* todo_flags_finish */ +}; + +class pass_omp_target_link : public gimple_opt_pass +{ +public: + pass_omp_target_link (gcc::context *ctxt) + : gimple_opt_pass (pass_data_omp_target_link, ctxt) + {} + + /* opt_pass methods: */ + virtual bool gate (function *fun) + { +#ifdef ACCEL_COMPILER + tree attrs = DECL_ATTRIBUTES (fun->decl); + return lookup_attribute ("omp declare target", attrs) + || lookup_attribute ("omp target entrypoint", attrs); +#else + (void) fun; + return false; +#endif + } + + virtual unsigned execute (function *); +}; + +/* Callback for walk_gimple_stmt used to scan for link var operands. */ + +static tree +find_link_var_op (tree *tp, int *walk_subtrees, void *) +{ + tree t = *tp; + + if (TREE_CODE (t) == VAR_DECL && DECL_HAS_VALUE_EXPR_P (t) + && lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (t))) + { + *walk_subtrees = 0; + return t; + } + + return NULL_TREE; +} + +unsigned +pass_omp_target_link::execute (function *fun) +{ + basic_block bb; + FOR_EACH_BB_FN (bb, fun) + { + gimple_stmt_iterator gsi; + for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi)) + if (walk_gimple_stmt (&gsi, NULL, find_link_var_op, NULL)) + gimple_regimplify_operands (gsi_stmt (gsi), &gsi); + } + + return 0; +} + +} // anon namespace + +gimple_opt_pass * +make_pass_omp_target_link (gcc::context *ctxt) +{ + return new pass_omp_target_link (ctxt); +} + #include "gt-omp-low.h" diff --git a/gcc/passes.def b/gcc/passes.def index 43ce3d5f1bf..c72b38ba1e5 100644 --- a/gcc/passes.def +++ b/gcc/passes.def @@ -170,6 +170,7 @@ along with GCC; see the file COPYING3. If not see NEXT_PASS (pass_fixup_cfg); NEXT_PASS (pass_lower_eh_dispatch); NEXT_PASS (pass_oacc_device_lower); + NEXT_PASS (pass_omp_target_link); NEXT_PASS (pass_all_optimizations); PUSH_INSERT_PASSES_WITHIN (pass_all_optimizations) NEXT_PASS (pass_remove_cgraph_callee_edges); diff --git a/gcc/tree-pass.h b/gcc/tree-pass.h index e1cbce923e4..a13a8653f23 100644 --- a/gcc/tree-pass.h +++ b/gcc/tree-pass.h @@ -417,6 +417,7 @@ extern gimple_opt_pass *make_pass_lower_omp (gcc::context *ctxt); extern gimple_opt_pass *make_pass_diagnose_omp_blocks (gcc::context *ctxt); extern gimple_opt_pass *make_pass_expand_omp (gcc::context *ctxt); extern gimple_opt_pass *make_pass_expand_omp_ssa (gcc::context *ctxt); +extern gimple_opt_pass *make_pass_omp_target_link (gcc::context *ctxt); extern gimple_opt_pass *make_pass_oacc_device_lower (gcc::context *ctxt); extern gimple_opt_pass *make_pass_object_sizes (gcc::context *ctxt); extern gimple_opt_pass *make_pass_strlen (gcc::context *ctxt); diff --git a/gcc/varpool.c b/gcc/varpool.c index 5e4fcbf2a4a..d0101a10818 100644 --- a/gcc/varpool.c +++ b/gcc/varpool.c @@ -748,6 +748,13 @@ symbol_table::output_variables (void) /* Handled in output_in_order. */ if (node->no_reorder) continue; +#ifdef ACCEL_COMPILER + /* Do not assemble "omp declare target link" vars. */ + if (DECL_HAS_VALUE_EXPR_P (node->decl) + && lookup_attribute ("omp declare target link", + DECL_ATTRIBUTES (node->decl))) + continue; +#endif if (node->assemble_decl ()) changed = true; } diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 87459270a5e..9315d8b83b2 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,21 @@ +2015-12-15 Ilya Verbin + + * libgomp.h (REFCOUNT_LINK): Define. + (struct splay_tree_key_s): Add link_key. + * target.c (gomp_map_vars): Treat REFCOUNT_LINK objects as not mapped. + Replace target address of the pointer with target address of newly + mapped object in the splay tree. Set link pointer on target to the + device address of the mapped object. + (gomp_unmap_vars): Restore target address of the pointer in the splay + tree for REFCOUNT_LINK objects after unmapping. + (gomp_load_image_to_device): Set refcount to REFCOUNT_LINK for "omp + declare target link" objects. + (gomp_unload_image_from_device): Replace j with i. Force unmap of all + "omp declare target link" objects, which were mapped for the image. + (gomp_exit_data): Restore target address of the pointer in the splay + tree for REFCOUNT_LINK objects after unmapping. + * testsuite/libgomp.c/target-link-1.c: New file. + 2015-12-14 Ilya Verbin * libgomp.h (gomp_device_state): New enum. diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 9d9949ff16a..73aa513b47c 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -817,6 +817,9 @@ struct target_mem_desc { /* Special value for refcount - infinity. */ #define REFCOUNT_INFINITY (~(uintptr_t) 0) +/* Special value for refcount - tgt_offset contains target address of the + artificial pointer to "omp declare target link" object. */ +#define REFCOUNT_LINK (~(uintptr_t) 1) struct splay_tree_key_s { /* Address of the host object. */ @@ -831,6 +834,8 @@ struct splay_tree_key_s { uintptr_t refcount; /* Asynchronous reference count. */ uintptr_t async_refcount; + /* Pointer to the original mapping of "omp declare target link" object. */ + splay_tree_key link_key; }; /* The comparison function. */ diff --git a/libgomp/target.c b/libgomp/target.c index 932b1761f4e..1ab30f7ca25 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -464,7 +464,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, } else n = splay_tree_lookup (mem_map, &cur_node); - if (n) + if (n && n->refcount != REFCOUNT_LINK) gomp_map_vars_existing (devicep, n, &cur_node, &tgt->list[i], kind & typemask); else @@ -628,11 +628,19 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, else k->host_end = k->host_start + sizeof (void *); splay_tree_key n = splay_tree_lookup (mem_map, k); - if (n) + if (n && n->refcount != REFCOUNT_LINK) gomp_map_vars_existing (devicep, n, k, &tgt->list[i], kind & typemask); else { + k->link_key = NULL; + if (n && n->refcount == REFCOUNT_LINK) + { + /* Replace target address of the pointer with target address + of mapped object in the splay tree. */ + splay_tree_remove (mem_map, n); + k->link_key = n; + } size_t align = (size_t) 1 << (kind >> rshift); tgt->list[i].key = k; k->tgt = tgt; @@ -752,6 +760,16 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, gomp_fatal ("%s: unhandled kind 0x%.2x", __FUNCTION__, kind); } + + if (k->link_key) + { + /* Set link pointer on target to the device address of the + mapped object. */ + void *tgt_addr = (void *) (tgt->tgt_start + k->tgt_offset); + devicep->host2dev_func (devicep->target_id, + (void *) n->tgt_offset, + &tgt_addr, sizeof (void *)); + } array++; } } @@ -884,6 +902,9 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom) if (do_unmap) { splay_tree_remove (&devicep->mem_map, k); + if (k->link_key) + splay_tree_insert (&devicep->mem_map, + (splay_tree_node) k->link_key); if (k->tgt->refcount > 1) k->tgt->refcount--; else @@ -1020,31 +1041,40 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version, k->tgt_offset = target_table[i].start; k->refcount = REFCOUNT_INFINITY; k->async_refcount = 0; + k->link_key = NULL; array->left = NULL; array->right = NULL; splay_tree_insert (&devicep->mem_map, array); array++; } + /* Most significant bit of the size in host and target tables marks + "omp declare target link" variables. */ + const uintptr_t link_bit = 1ULL << (sizeof (uintptr_t) * __CHAR_BIT__ - 1); + const uintptr_t size_mask = ~link_bit; + for (i = 0; i < num_vars; i++) { struct addr_pair *target_var = &target_table[num_funcs + i]; - if (target_var->end - target_var->start - != (uintptr_t) host_var_table[i * 2 + 1]) + uintptr_t target_size = target_var->end - target_var->start; + + if ((uintptr_t) host_var_table[i * 2 + 1] != target_size) { gomp_mutex_unlock (&devicep->lock); if (is_register_lock) gomp_mutex_unlock (®ister_lock); - gomp_fatal ("Can't map target variables (size mismatch)"); + gomp_fatal ("Cannot map target variables (size mismatch)"); } splay_tree_key k = &array->key; k->host_start = (uintptr_t) host_var_table[i * 2]; - k->host_end = k->host_start + (uintptr_t) host_var_table[i * 2 + 1]; + k->host_end + = k->host_start + (size_mask & (uintptr_t) host_var_table[i * 2 + 1]); k->tgt = tgt; k->tgt_offset = target_var->start; - k->refcount = REFCOUNT_INFINITY; + k->refcount = target_size & link_bit ? REFCOUNT_LINK : REFCOUNT_INFINITY; k->async_refcount = 0; + k->link_key = NULL; array->left = NULL; array->right = NULL; splay_tree_insert (&devicep->mem_map, array); @@ -1072,7 +1102,6 @@ gomp_unload_image_from_device (struct gomp_device_descr *devicep, int num_funcs = host_funcs_end - host_func_table; int num_vars = (host_vars_end - host_var_table) / 2; - unsigned j; struct splay_tree_key_s k; splay_tree_key node = NULL; @@ -1088,21 +1117,46 @@ gomp_unload_image_from_device (struct gomp_device_descr *devicep, devicep->unload_image_func (devicep->target_id, version, target_data); /* Remove mappings from splay tree. */ - for (j = 0; j < num_funcs; j++) + int i; + for (i = 0; i < num_funcs; i++) { - k.host_start = (uintptr_t) host_func_table[j]; + k.host_start = (uintptr_t) host_func_table[i]; k.host_end = k.host_start + 1; splay_tree_remove (&devicep->mem_map, &k); } - for (j = 0; j < num_vars; j++) + /* Most significant bit of the size in host and target tables marks + "omp declare target link" variables. */ + const uintptr_t link_bit = 1ULL << (sizeof (uintptr_t) * __CHAR_BIT__ - 1); + const uintptr_t size_mask = ~link_bit; + bool is_tgt_unmapped = false; + + for (i = 0; i < num_vars; i++) { - k.host_start = (uintptr_t) host_var_table[j * 2]; - k.host_end = k.host_start + (uintptr_t) host_var_table[j * 2 + 1]; - splay_tree_remove (&devicep->mem_map, &k); + k.host_start = (uintptr_t) host_var_table[i * 2]; + k.host_end + = k.host_start + (size_mask & (uintptr_t) host_var_table[i * 2 + 1]); + + if (!(link_bit & (uintptr_t) host_var_table[i * 2 + 1])) + splay_tree_remove (&devicep->mem_map, &k); + else + { + splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &k); + splay_tree_remove (&devicep->mem_map, n); + if (n->link_key) + { + if (n->tgt->refcount > 1) + n->tgt->refcount--; + else + { + is_tgt_unmapped = true; + gomp_unmap_tgt (n->tgt); + } + } + } } - if (node) + if (node && !is_tgt_unmapped) { free (node->tgt); free (node); @@ -1658,6 +1712,9 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum, if (k->refcount == 0) { splay_tree_remove (&devicep->mem_map, k); + if (k->link_key) + splay_tree_insert (&devicep->mem_map, + (splay_tree_node) k->link_key); if (k->tgt->refcount > 1) k->tgt->refcount--; else diff --git a/libgomp/testsuite/libgomp.c/target-link-1.c b/libgomp/testsuite/libgomp.c/target-link-1.c new file mode 100644 index 00000000000..681677cc2aa --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-link-1.c @@ -0,0 +1,63 @@ +struct S { int s, t; }; + +int a = 1, b = 1; +double c[27]; +struct S d = { 8888, 8888 }; +#pragma omp declare target link (a) to (b) link (c, d) + +int +foo (void) +{ + return a++ + b++; +} + +int +bar (int n) +{ + int *p1 = &a; + int *p2 = &b; + c[n] += 2.0; + d.s -= 2; + d.t -= 2; + return *p1 + *p2 + d.s + d.t; +} + +#pragma omp declare target (foo, bar) + +int +main () +{ + a = b = 2; + d.s = 17; + d.t = 18; + + int res, n = 10; + #pragma omp target map (to: a, b, c, d) map (from: res) + { + res = foo () + foo (); + c[n] = 3.0; + res += bar (n); + } + + int shared_mem = 0; + #pragma omp target map (alloc: shared_mem) + shared_mem = 1; + + if ((shared_mem && res != (2 + 2) + (3 + 3) + (4 + 4 + 15 + 16)) + || (!shared_mem && res != (2 + 1) + (3 + 2) + (4 + 3 + 15 + 16))) + __builtin_abort (); + + #pragma omp target enter data map (to: c) + #pragma omp target update from (c) + res = (int) (c[n] + 0.5); + if ((shared_mem && res != 5) || (!shared_mem && res != 0)) + __builtin_abort (); + + #pragma omp target map (to: a, b) map (from: res) + res = foo (); + + if ((shared_mem && res != 4 + 4) || (!shared_mem && res != 2 + 3)) + __builtin_abort (); + + return 0; +}