https://gcc.gnu.org/g:7a8f98b48104ccf10c6dceccc51b70de69288eaa
commit r17-2003-g7a8f98b48104ccf10c6dceccc51b70de69288eaa Author: Chung-Lin Tang <[email protected]> Date: Tue Jun 30 09:15:22 2026 +0000 OpenMP: support for uses_allocators clause This patch completes the implementation of uses_allocators. A new map kind for libgomp GOMP_MAP_USES_ALLOCATORS has been implemented, which completes the required allocator initialization inside libgomp on the host side for the accel target. This bypasses the various problems of initializing on the accel target. To this end, the libgomp omp_allocator_data format must be strictly the same between host/accelerator, which is not a hard to satisfy requirement. New configure time checks have been added in this patch to enforce this. There is now however a duplication of the 'memspace_validate' logic, which before was only present in libgomp/config/<accel_tgt>/allocator.c, now also is duplicated in libgomp/plugin/plugin-<accel_tgt>.c, as a plugin hook to allow the host to use to validate for the accel target. Reviewed-by: Tobias Burnus <[email protected]> gcc/c/ChangeLog: * c-typeck.cc (c_finish_omp_clauses): Fix comments. gcc/fortran/ChangeLog: * trans-array.cc (gfc_conv_array_initializer): Always set PURPOSE when building constructor for get_initialized_tmp_var, adjust test to use integer_zerop instead of != NULL_TREE. * trans-openmp.cc (gfc_trans_omp_clauses): Handle uses_allocators. gcc/ChangeLog: * gimplify.cc (enum gimplify_omp_var_data): Add GOVD_USES_ALLOCATORS_ALLOCATOR. (gimplify_bind_expr): Diagnose missing uses_allocators clause. (omp_get_attachment): Add GOMP_MAP_USES_ALLOCATORS cases. (omp_group_base): Likewise. (gimplify_scan_omp_clauses): Handle uses_allocators. (gimplify_adjust_omp_clauses_1): Handle GOVD_USES_ALLOCATORS_ALLOCATOR. * omp-low.cc (scan_sharing_clauses): Check if ALLOCATE clause in target region has containing USES_ALLOCATORS clause. Add OMP_CLAUSE_USES_ALLOCATORS cases. (lower_omp_target): Add OMP_CLAUSE_USES_ALLOCATORS cases. * tree-pretty-print.cc (dump_omp_clause): Handle GOMP_MAP_USES_ALLOCATORS. gcc/testsuite/ChangeLog: * c-c++-common/gomp/uses_allocators-1.c: Adjust test, and move to ../../libgomp/testsuite/libgomp.c-c++-common/uses_allocators-1.c * c-c++-common/gomp/uses_allocators-2.c: Move to ../../libgomp/testsuite/libgomp.c-c++-common/uses_allocators-2.c * c-c++-common/gomp/uses_allocators-4.c: Adjust test, and move to ../../libgomp/testsuite/libgomp.c-c++-common/uses_allocators-4.c * c-c++-common/gomp/uses_allocators-7.c: Move to ../../libgomp/testsuite/libgomp.c-c++-common/uses_allocators-7.c * c-c++-common/gomp/uses_allocators-8.c: Move to ../../libgomp/testsuite/libgomp.c-c++-common/uses_allocators-8.c * gcc.dg/gomp/deprecate-2.c: Adjust test. * gfortran.dg/gomp/allocate-1.f90: Likewise. * gfortran.dg/gomp/scope-6.f90: Likewise. * g++.dg/gomp/uses_allocators-1.C: Move to ../../libgomp/testsuite/libgomp.c++/uses_allocators-1.C. include/ChangeLog: * gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_USES_ALLOCATORS. libgomp/ChangeLog: * allocator.c (struct omp_allocator_data): Always enable 'memkind' field for configurations with offloading. Adjust comments assuming HAVE_SYNC_BUILTINS for all offload host/accelerators. (gomp_init_allocator_data): New function, created from omp_init_allocator logic. (omp_init_allocator): Adjust to use gomp_init_allocator_data. (gomp_map_omp_init_allocator): New function to call gomp_init_allocator_data for uses_allocators clauses. * config/gcn/allocator.c (gcn_memspace_validate): Add comments. * config/nvptx/allocator.c (nvptx_memspace_alloc): Add PTX >= v4.1 preprocessor symbol check. (nvptx_memspace_calloc): Likewise. (nvptx_memspace_free): Likewise. (nvptx_memspace_realloc): Likewise. (nvptx_memspace_validate): Add comments. * configure.ac: Require offload hosts and accelerator targets to support __sync_* builtins. * configure: Regenerate. * libgomp-plugin.h (GOMP_OFFLOAD_memspace_validate): New declaration. * libgomp.h (gomp_omp_allocator_data_size): New declaration. (memspace_validate_func): New plugin hook in struct gomp_device_descr. (gomp_map_omp_init_allocator): New declaration. * oacc-host.c (host_dispatch): Init .memspace_validate_func = NULL. * plugin/plugin-gcn.c (GOMP_OFFLOAD_memspace_validate): New function. * plugin/plugin-nvptx.c (GOMP_OFFLOAD_memspace_validate): New function. * target.c (gomp_map_vars_internal): Implement GOMP_MAP_USES_ALLOCATORS map kind. (calculate_firstprivate_requirements): Likewise. (copy_firstprivate_data): Likewise. (gomp_load_plugin_for_device): Add DLSYM_OPT of memspace_validate_func hook. * testsuite/libgomp.c++/uses_allocators-1.C: Moved from ../gcc/testsuite/g++.dg/gomp/uses_allocators-1.C. * testsuite/libgomp.c-c++-common/uses_allocators-1.c: Moved from ../gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c. * testsuite/libgomp.c-c++-common/uses_allocators-2.c: Moved from ../gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c. * testsuite/libgomp.c-c++-common/uses_allocators-4.c: Moved from ../gcc/testsuite/c-c++-common/gomp/uses_allocators-4.c. * testsuite/libgomp.c-c++-common/uses_allocators-7.c: Moved from ../gcc/testsuite/c-c++-common/gomp/uses_allocators-7.c. * testsuite/libgomp.c-c++-common/uses_allocators-8.c: Moved from ../gcc/testsuite/c-c++-common/gomp/uses_allocators-8.c. * testsuite/libgomp.c-c++-common/uses_allocators-3.c: New test. * testsuite/libgomp.c-c++-common/uses_allocators-5.c: New test. * testsuite/libgomp.c-c++-common/uses_allocators-6.c: New test. * testsuite/libgomp.c-c++-common/uses_allocators-9.c: New test. * testsuite/libgomp.fortran/uses_allocators-1.f90: New test. * testsuite/libgomp.fortran/uses_allocators-2.f90: Rename from uses_allocators_1.f90. * testsuite/libgomp.fortran/uses_allocators-3.f90: Rename from uses_allocators_2.f90. * testsuite/libgomp.fortran/uses_allocators-4.f90: New test. * testsuite/libgomp.fortran/uses_allocators-5.f90: New test. * testsuite/libgomp.fortran/uses_allocators-6.f90: New test. * testsuite/libgomp.fortran/uses_allocators-7.f90: Adjust test. * testsuite/libgomp.fortran/uses_allocators-8.f90: New test. * testsuite/libgomp.fortran/uses_allocators_1.f90: Renamed to uses_allocators-2.f90 and deleted. * testsuite/libgomp.fortran/uses_allocators_2.f90: Renamed to uses_allocators-3.f90 and deleted. * testsuite/libgomp.c++/c++.exp (check_effective_target_c): Add. (check_effective_target_c++): Add. * testsuite/libgomp.c/c.exp (check_effective_target_c): Add. (check_effective_target_c++): Add. Diff: --- gcc/c/c-typeck.cc | 2 +- gcc/fortran/trans-array.cc | 7 +- gcc/fortran/trans-openmp.cc | 43 ++++- gcc/gimplify.cc | 100 ++++++++++- gcc/omp-low.cc | 121 ++++++++++++++ .../c-c++-common/gomp/uses_allocators-7.c | 83 ---------- .../c-c++-common/gomp/uses_allocators-8.c | 61 ------- gcc/testsuite/gcc.dg/gomp/deprecate-2.c | 3 - gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 | 7 +- gcc/testsuite/gfortran.dg/gomp/scope-6.f90 | 2 +- gcc/tree-pretty-print.cc | 3 + include/gomp-constants.h | 5 + libgomp/allocator.c | 184 +++++++++++++++------ libgomp/config/gcn/allocator.c | 4 + libgomp/config/nvptx/allocator.c | 31 +++- libgomp/configure | 19 +++ libgomp/configure.ac | 18 ++ libgomp/libgomp-plugin.h | 1 + libgomp/libgomp.h | 9 + libgomp/oacc-host.c | 1 + libgomp/plugin/plugin-gcn.c | 16 ++ libgomp/plugin/plugin-nvptx.c | 18 ++ libgomp/target.c | 56 ++++++- libgomp/testsuite/libgomp.c++/c++.exp | 9 + .../testsuite/libgomp.c++}/uses_allocators-1.C | 10 +- .../libgomp.c-c++-common}/uses_allocators-1.c | 25 +-- .../libgomp.c-c++-common}/uses_allocators-2.c | 19 +-- .../libgomp.c-c++-common/uses_allocators-3.c | 50 ++++++ .../libgomp.c-c++-common}/uses_allocators-4.c | 88 +--------- .../libgomp.c-c++-common/uses_allocators-5.c | 34 ++++ .../libgomp.c-c++-common/uses_allocators-6.c | 53 ++++++ .../libgomp.c-c++-common/uses_allocators-7.c | 35 ++++ .../libgomp.c-c++-common/uses_allocators-8.c | 16 ++ .../libgomp.c-c++-common/uses_allocators-9.c | 44 +++++ libgomp/testsuite/libgomp.c/c.exp | 8 + .../libgomp.fortran/uses_allocators-1.f90 | 18 ++ ...uses_allocators_1.f90 => uses_allocators-2.f90} | 0 ...uses_allocators_2.f90 => uses_allocators-3.f90} | 22 ++- .../libgomp.fortran/uses_allocators-4.f90 | 59 +++++++ .../libgomp.fortran/uses_allocators-5.f90 | 54 ++++++ .../libgomp.fortran/uses_allocators-6.f90 | 14 ++ .../libgomp.fortran/uses_allocators-7.f90 | 21 +-- .../libgomp.fortran/uses_allocators-8.f90 | 51 ++++++ 43 files changed, 1042 insertions(+), 382 deletions(-) diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index c8127cc1f053..643035b50a08 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -17419,7 +17419,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else if (TREE_CODE (t) == CONST_DECL) { /* omp_null_allocator is ignored and for predefined allocators, - not special handling is required; thus, remove them removed. */ + no special handling is required; thus, mark them removed. */ remove = true; if (OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) diff --git a/gcc/fortran/trans-array.cc b/gcc/fortran/trans-array.cc index ebd04c53b9d2..6ff5ac589fed 100644 --- a/gcc/fortran/trans-array.cc +++ b/gcc/fortran/trans-array.cc @@ -7243,10 +7243,7 @@ gfc_conv_array_initializer (tree type, gfc_expr * expr) &expr->where, flag_max_array_constructor); return NULL_TREE; } - if (mpz_cmp_si (c->offset, 0) != 0) - index = gfc_conv_mpz_to_tree (c->offset, gfc_index_integer_kind); - else - index = NULL_TREE; + index = gfc_conv_mpz_to_tree (c->offset, gfc_index_integer_kind); if (mpz_cmp_si (c->repeat, 1) > 0) { @@ -7317,7 +7314,7 @@ gfc_conv_array_initializer (tree type, gfc_expr * expr) CONSTRUCTOR_APPEND_ELT (v, index, se.expr); else { - if (index != NULL_TREE) + if (!integer_zerop (index)) CONSTRUCTOR_APPEND_ELT (v, index, se.expr); CONSTRUCTOR_APPEND_ELT (v, range, se.expr); } diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 8d5842685a55..36ce4f038e95 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -4061,7 +4061,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, gfc_init_se (&se, NULL); gfc_conv_expr (&se, n->u2.allocator); gfc_add_block_to_block (block, &se.pre); - allocator_ = gfc_evaluate_now (se.expr, block); + t = se.expr; + if (DECL_P (t) && se.post.head == NULL_TREE) + allocator_ = (POINTER_TYPE_P (TREE_TYPE (t)) + ? build_fold_indirect_ref (t): t); + else + allocator_ = gfc_evaluate_now (t, block); gfc_add_block_to_block (block, &se.post); } OMP_CLAUSE_ALLOCATE_ALLOCATOR (node) = allocator_; @@ -5471,14 +5476,36 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, } break; case OMP_LIST_USES_ALLOCATORS: - /* Ignore omp_null_allocator and pre-defined allocators as no - special treatment is needed. */ for (; n != NULL; n = n->next) - if (n->sym->attr.flavor == FL_VARIABLE) - break; - if (n != NULL) - sorry_at (input_location, "%<uses_allocators%> clause with traits " - "and memory spaces"); + { + if (!n->sym->attr.referenced) + continue; + tree node = build_omp_clause (input_location, + OMP_CLAUSE_USES_ALLOCATORS); + tree t; + if (n->sym->attr.flavor == FL_VARIABLE) + t = gfc_get_symbol_decl (n->sym); + else + { + t = gfc_conv_mpz_to_tree (n->sym->value->value.integer, + n->sym->ts.kind); + t = fold_convert (ptr_type_node, t); + } + OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR(node) = t; + if (n->u.memspace_sym) + { + n->u.memspace_sym->attr.referenced = true; + OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (node) + = gfc_get_symbol_decl (n->u.memspace_sym); + } + if (n->u2.traits_sym) + { + n->u2.traits_sym->attr.referenced = true; + OMP_CLAUSE_USES_ALLOCATORS_TRAITS (node) + = gfc_get_symbol_decl (n->u2.traits_sym); + } + omp_clauses = gfc_trans_add_clause (node, omp_clauses); + } break; default: break; diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 8f64979571b7..6c5d182ffab2 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -185,6 +185,9 @@ enum gimplify_omp_var_data /* Flag for GOVD_FIRSTPRIVATE: OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT. */ GOVD_FIRSTPRIVATE_IMPLICIT = 0x4000000, + /* Flag to indicate this an allocator for the uses_allocators clause. */ + GOVD_USES_ALLOCATORS_ALLOCATOR = 0x8000000, + GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR | GOVD_LOCAL) @@ -1291,18 +1294,46 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p) dynamic_allocators clause is present in the same compilation unit. */ bool missing_dyn_alloc = false; - if (alloc == NULL_TREE - && ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) - == 0)) + if ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0) { /* This comes too early for omp_discover_declare_target..., but should at least catch the most common cases. */ missing_dyn_alloc - = cgraph_node::get (current_function_decl)->offloadable; + = (alloc == NULL_TREE + && cgraph_node::get (current_function_decl)->offloadable); for (struct gimplify_omp_ctx *ctx2 = ctx; ctx2 && !missing_dyn_alloc; ctx2 = ctx2->outer_context) if (ctx2->code == OMP_TARGET) - missing_dyn_alloc = true; + { + if (alloc == NULL_TREE) + missing_dyn_alloc = true; + else if (TREE_CODE (alloc) != INTEGER_CST) + { + tree alloc2 = alloc; + if (TREE_CODE (alloc2) == MEM_REF + || TREE_CODE (alloc2) == INDIRECT_REF) + alloc2 = TREE_OPERAND (alloc2, 0); + tree c2; + for (c2 = ctx2->clauses; c2; + c2 = OMP_CLAUSE_CHAIN (c2)) + if (OMP_CLAUSE_CODE (c2) + == OMP_CLAUSE_USES_ALLOCATORS) + { + tree t2 + = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c2); + if (operand_equal_p (alloc2, t2)) + break; + } + if (c2 == NULL_TREE) + error_at (EXPR_LOC_OR_LOC ( + alloc, DECL_SOURCE_LOCATION (t)), + "%qE in %<allocator%> clause inside a " + "target region must be specified in an " + "%<uses_allocators%> clause on the " + "%<target%> directive", alloc2); + } + break; + } } if (missing_dyn_alloc) error_at (DECL_SOURCE_LOCATION (t), @@ -10866,6 +10897,7 @@ omp_get_attachment (omp_mapping_group *grp) case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + case GOMP_MAP_USES_ALLOCATORS: return NULL_TREE; case GOMP_MAP_ATTACH_DETACH: @@ -10910,6 +10942,7 @@ omp_get_attachment (omp_mapping_group *grp) case GOMP_MAP_FIRSTPRIVATE_INT: case GOMP_MAP_USE_DEVICE_PTR: case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + case GOMP_MAP_USES_ALLOCATORS: return NULL_TREE; default: @@ -11198,6 +11231,7 @@ omp_group_base (omp_mapping_group *grp, unsigned int *chained, case GOMP_MAP_FIRSTPRIVATE_INT: case GOMP_MAP_USE_DEVICE_PTR: case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + case GOMP_MAP_USES_ALLOCATORS: return NULL_TREE; case GOMP_MAP_FIRSTPRIVATE_POINTER: @@ -14824,8 +14858,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; case OMP_CLAUSE_USES_ALLOCATORS: - sorry_at (OMP_CLAUSE_LOCATION (c), "%<uses_allocators%> clause"); - remove = 1; + if (TREE_CODE (OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c)) + != INTEGER_CST) + { + decl = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c); + omp_add_variable (ctx, decl, + GOVD_SEEN | GOVD_USES_ALLOCATORS_ALLOCATOR); + } + else + remove = true; break; case OMP_CLAUSE_ORDERED: @@ -14984,6 +15025,49 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, remove = true; break; } + if ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0 + && OMP_CLAUSE_ALLOCATE_ALLOCATOR (c) + && TREE_CODE (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)) != INTEGER_CST) + { + tree allocator = OMP_CLAUSE_ALLOCATE_ALLOCATOR (c); + tree clauses = NULL_TREE; + + /* Get clause list of the nearest enclosing target construct. */ + if (ctx->code == OMP_TARGET) + clauses = *orig_list_p; + else + { + struct gimplify_omp_ctx *tctx = ctx->outer_context; + while (tctx && tctx->code != OMP_TARGET) + tctx = tctx->outer_context; + if (tctx) + clauses = tctx->clauses; + } + + if (clauses) + { + tree uc; + if (TREE_CODE (allocator) == MEM_REF + || TREE_CODE (allocator) == INDIRECT_REF) + allocator = TREE_OPERAND (allocator, 0); + for (uc = clauses; uc; uc = OMP_CLAUSE_CHAIN (uc)) + if (OMP_CLAUSE_CODE (uc) == OMP_CLAUSE_USES_ALLOCATORS) + { + tree uc_allocator + = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (uc); + if (operand_equal_p (allocator, uc_allocator)) + break; + } + if (uc == NULL_TREE) + { + error_at (OMP_CLAUSE_LOCATION (c), "allocator %qE " + "requires %<uses_allocators(%E)%> clause in " + "target region", allocator, allocator); + remove = true; + break; + } + } + } if (gimplify_expr (&OMP_CLAUSE_ALLOCATE_ALLOCATOR (c), pre_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR) { @@ -15302,6 +15386,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) code = OMP_CLAUSE__CONDTEMP_; gimple_add_tmp_var (decl); } + else if (flags & GOVD_USES_ALLOCATORS_ALLOCATOR) + return 0; else gcc_unreachable (); diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index dbdbc4079f4d..ad5b2225d279 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -1193,6 +1193,36 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && omp_maybe_offloaded_ctx (ctx)) error_at (OMP_CLAUSE_LOCATION (c), "%<allocate%> clause must" " specify an allocator here"); + if ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0 + && OMP_CLAUSE_ALLOCATE_ALLOCATOR (c) != NULL_TREE + && DECL_P (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)) + && !DECL_ARTIFICIAL (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c))) + { + tree alloc2 = OMP_CLAUSE_ALLOCATE_ALLOCATOR (c); + if (TREE_CODE (alloc2) == MEM_REF + || TREE_CODE (alloc2) == INDIRECT_REF) + alloc2 = TREE_OPERAND (alloc2, 0); + omp_context *ctx2 = ctx; + for (; ctx2; ctx2 = ctx2->outer) + if (is_gimple_omp_offloaded (ctx2->stmt)) + break; + if (ctx2 != NULL) + { + tree c2 = gimple_omp_target_clauses (ctx2->stmt); + for (; c2; c2 = OMP_CLAUSE_CHAIN (c2)) + if (OMP_CLAUSE_CODE (c2) == OMP_CLAUSE_USES_ALLOCATORS + && operand_equal_p ( + alloc2, OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c2))) + break; + if (c2 == NULL_TREE) + error_at (EXPR_LOC_OR_LOC (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c), + OMP_CLAUSE_LOCATION (c)), + "allocator %qE in %<allocate%> clause inside a " + "target region must be specified in an " + "%<uses_allocators%> clause on the %<target%> " + "directive", alloc2); + } + } if (ctx->allocate_map == NULL) ctx->allocate_map = new hash_map<tree, tree>; tree val = integer_zero_node; @@ -1797,6 +1827,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_DEVICE_TYPE: break; + case OMP_CLAUSE_USES_ALLOCATORS: + decl = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c); + gcc_assert (DECL_P (decl)); + gcc_assert (is_gimple_omp_offloaded (ctx->stmt)); + install_var_field (decl, false, 3, ctx); + install_var_local (decl, ctx); + break; + case OMP_CLAUSE_ALIGNED: decl = OMP_CLAUSE_DECL (c); if (is_global_var (decl) @@ -2022,6 +2060,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_USE: case OMP_CLAUSE_DESTROY: case OMP_CLAUSE_DEVICE_TYPE: + case OMP_CLAUSE_USES_ALLOCATORS: break; case OMP_CLAUSE__CACHE_: @@ -13233,6 +13272,14 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) sorry_at (OMP_CLAUSE_LOCATION (c), "only the %<device_type(any)%> is supported"); break; + case OMP_CLAUSE_USES_ALLOCATORS: + allocator = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c); + tree new_allocator = lookup_decl (allocator, ctx); + x = build_receiver_ref (allocator, false, ctx); + SET_DECL_VALUE_EXPR (new_allocator, x); + DECL_HAS_VALUE_EXPR_P (new_allocator) = 1; + map_cnt++; + break; } if (offloaded) @@ -13966,6 +14013,80 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) CONSTRUCTOR_APPEND_ELT (vkind, purpose, build_int_cstu (tkind_type, tkind)); break; + + case OMP_CLAUSE_USES_ALLOCATORS: + tree allocator = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c); + tree memspace = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c); + tree traits = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c); + + tree ntraits, traits_var; + if (traits == NULL_TREE) + { + ntraits = integer_zero_node; + traits_var = null_pointer_node; + } + else if (DECL_INITIAL (traits)) + { + location_t loc = OMP_CLAUSE_LOCATION (c); + ntraits = array_type_nelts_top (TREE_TYPE (traits)); + tree t = DECL_INITIAL (traits); + t = get_initialized_tmp_var (t, &ilist, NULL); + traits_var = build_fold_addr_expr_loc (loc, t); + } + else + { + /* This happens for VLAs, which probably aren't useful + because they can't be const initialized in the same + scope.... is there something else? */ + location_t loc = OMP_CLAUSE_LOCATION (c); + gcc_assert (TREE_CODE (TREE_TYPE (traits)) == ARRAY_TYPE); + ntraits = array_type_nelts_top (TREE_TYPE (traits)); + traits_var = build_fold_addr_expr_loc (loc, traits); + } + + if (memspace == NULL_TREE) + memspace = build_int_cst (pointer_sized_int_node, 0); + else + memspace = fold_convert (pointer_sized_int_node, memspace); + + tree arr_type + = build_array_type_nelts (pointer_sized_int_node, 3); + tree uses_allocators_descr + = create_tmp_var (arr_type, "uses_allocator_descr"); + tree ua_descr[3]; + for (int i = 0; i < 3; i++) + ua_descr[i] = build4 (ARRAY_REF, pointer_sized_int_node, + uses_allocators_descr, + build_int_cst (size_type_node, i), + NULL_TREE, NULL_TREE); + gimplify_assign (ua_descr[0], + fold_convert (pointer_sized_int_node, memspace), + &ilist); + gimplify_assign (ua_descr[1], + fold_convert (pointer_sized_int_node, ntraits), + &ilist); + gimplify_assign (ua_descr[2], + fold_convert (pointer_sized_int_node, traits_var), + &ilist); + + x = build_sender_ref (allocator, ctx); + tree ptr = build_fold_addr_expr (uses_allocators_descr); + gimplify_assign (x, fold_convert (TREE_TYPE (x), ptr), &ilist); + + s = size_int (0); + purpose = size_int (map_idx++); + CONSTRUCTOR_APPEND_ELT (vsize, purpose, s); + tkind = GOMP_MAP_USES_ALLOCATORS; + gcc_checking_assert (tkind + < (HOST_WIDE_INT_C (1U) << talign_shift)); + talign = TYPE_ALIGN_UNIT (pointer_sized_int_node); + talign = ceil_log2 (talign); + tkind |= talign << talign_shift; + gcc_checking_assert (tkind + <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type))); + CONSTRUCTOR_APPEND_ELT (vkind, purpose, + build_int_cstu (tkind_type, tkind)); + break; } gcc_assert (map_idx == map_cnt); diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-7.c b/gcc/testsuite/c-c++-common/gomp/uses_allocators-7.c deleted file mode 100644 index ce43d70ecf25..000000000000 --- a/gcc/testsuite/c-c++-common/gomp/uses_allocators-7.c +++ /dev/null @@ -1,83 +0,0 @@ -// { dg-do compile } - -//#include <omp.h> - -typedef __UINTPTR_TYPE__ omp_uintptr_t; - -#if __cplusplus >= 201103L -# define __GOMP_UINTPTR_T_ENUM : omp_uintptr_t -#else -# define __GOMP_UINTPTR_T_ENUM -#endif - -typedef enum omp_memspace_handle_t __GOMP_UINTPTR_T_ENUM -{ - omp_default_mem_space = 0, - omp_large_cap_mem_space = 1, - omp_const_mem_space = 2, - omp_high_bw_mem_space = 3, - omp_low_lat_mem_space = 4, - ompx_gnu_managed_mem_space = 200, - __omp_memspace_handle_t_max__ = __UINTPTR_MAX__ -} omp_memspace_handle_t; - -typedef enum omp_allocator_handle_t __GOMP_UINTPTR_T_ENUM -{ - omp_null_allocator = 0, - omp_default_mem_alloc = 1, - omp_large_cap_mem_alloc = 2, - omp_const_mem_alloc = 3, - omp_high_bw_mem_alloc = 4, - omp_low_lat_mem_alloc = 5, - omp_cgroup_mem_alloc = 6, - omp_pteam_mem_alloc = 7, - omp_thread_mem_alloc = 8, - ompx_gnu_pinned_mem_alloc = 200, - ompx_gnu_managed_mem_alloc = 201, - __omp_allocator_handle_t_max__ = __UINTPTR_MAX__ -} omp_allocator_handle_t; - -typedef struct omp_alloctrait_t -{ -// omp_alloctrait_key_t key; -// omp_uintptr_t value; -} omp_alloctrait_t; - - -void f() -{ - omp_allocator_handle_t my; - struct t { - omp_allocator_handle_t a1; - } s; - -const omp_allocator_handle_t my3 = (omp_allocator_handle_t) 300; -omp_allocator_handle_t my4[1]; -const omp_alloctrait_t t[] = {}; - #pragma omp target uses_allocators(my, omp_default_mem_alloc, omp_null_allocator) // OK - ; - #pragma omp target uses_allocators(my) firstprivate(my) // { dg-error "'my' appears more than once in data clauses" } - ; - #pragma omp target private(my) uses_allocators(my) // { dg-error "'my' appears more than once in data clauses" } - ; - #pragma omp target uses_allocators(my3) - ; - #pragma omp target uses_allocators(s.a1) - // { dg-error "expected '\\)' before '\\.' token" "" { target *-*-* } .-1 } - // { dg-error "allocator 's' must be of 'omp_allocator_handle_t' type" "" { target *-*-* } .-2 } - ; - #pragma omp target uses_allocators(my4) - // { dg-error "allocator 'my4' must be of 'omp_allocator_handle_t' type" "" { target *-*-* } .-1 } - ; - #pragma omp target uses_allocators(my4[0]) - // { dg-error "expected '\\)' before '\\\[' token" "" { target *-*-* } .-1 } - // { dg-error "allocator 'my4' must be of 'omp_allocator_handle_t' type" "" { target *-*-* } .-2 } - ; - #pragma omp target uses_allocators(memspace(omp_default_mem_space) : my, my(t)) // { dg-error "legacy 'my\\(t\\)' traits syntax not allowed in 'uses_allocators' clause when using modifiers" } - ; -} - -// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 57 } -// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 61 } -// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 63 } -// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 76 } diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-8.c b/gcc/testsuite/c-c++-common/gomp/uses_allocators-8.c deleted file mode 100644 index 3ab87576c415..000000000000 --- a/gcc/testsuite/c-c++-common/gomp/uses_allocators-8.c +++ /dev/null @@ -1,61 +0,0 @@ -// { dg-do compile } -// { dg-additional-options "-fdump-tree-original" } - -//#include <omp.h> - -typedef __UINTPTR_TYPE__ omp_uintptr_t; - -#if __cplusplus >= 201103L -# define __GOMP_UINTPTR_T_ENUM : omp_uintptr_t -#else -# define __GOMP_UINTPTR_T_ENUM -#endif - -typedef enum omp_memspace_handle_t __GOMP_UINTPTR_T_ENUM -{ - omp_default_mem_space = 0, - omp_large_cap_mem_space = 1, - omp_const_mem_space = 2, - omp_high_bw_mem_space = 3, - omp_low_lat_mem_space = 4, - ompx_gnu_managed_mem_space = 200, - __omp_memspace_handle_t_max__ = __UINTPTR_MAX__ -} omp_memspace_handle_t; - -typedef enum omp_allocator_handle_t __GOMP_UINTPTR_T_ENUM -{ - omp_null_allocator = 0, - omp_default_mem_alloc = 1, - omp_large_cap_mem_alloc = 2, - omp_const_mem_alloc = 3, - omp_high_bw_mem_alloc = 4, - omp_low_lat_mem_alloc = 5, - omp_cgroup_mem_alloc = 6, - omp_pteam_mem_alloc = 7, - omp_thread_mem_alloc = 8, - ompx_gnu_pinned_mem_alloc = 200, - ompx_gnu_managed_mem_alloc = 201, - __omp_allocator_handle_t_max__ = __UINTPTR_MAX__ -} omp_allocator_handle_t; - -typedef struct omp_alloctrait_t -{ -// omp_alloctrait_key_t key; -// omp_uintptr_t value; -} omp_alloctrait_t; - - -void f() -{ - omp_allocator_handle_t my, my2, my3, my4; -const omp_alloctrait_t t[] = {}; -const omp_alloctrait_t t2[] = {}; - #pragma omp target uses_allocators(traits(t), memspace(omp_high_bw_mem_space) : my; omp_default_mem_alloc, omp_null_allocator; my2; traits(t2) : my3; memspace(omp_large_cap_mem_space) : my4) - ; -} - -// { dg-final { scan-tree-dump "#pragma omp target uses_allocators\\(memspace\\(1\\), traits\\(\\) : my4\\) uses_allocators\\(memspace\\(\\), traits\\(t2\\) : my3\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : my2\\) uses_allocators\\(memspace\\(3\\), traits\\(t\\) : my\\)" "original" { target c } } } -// { dg-final { scan-tree-dump "#pragma omp target uses_allocators\\(memspace\\(omp_large_cap_mem_space\\), traits\\(\\) : my4\\) uses_allocators\\(memspace\\(\\), traits\\(t2\\) : my3\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : my2\\) uses_allocators\\(memspace\\(omp_high_bw_mem_space\\), traits\\(t\\) : my\\)" "original" { target c++ } } } - - -// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 53 } diff --git a/gcc/testsuite/gcc.dg/gomp/deprecate-2.c b/gcc/testsuite/gcc.dg/gomp/deprecate-2.c index 46d286141ef2..b7689dd1554d 100644 --- a/gcc/testsuite/gcc.dg/gomp/deprecate-2.c +++ b/gcc/testsuite/gcc.dg/gomp/deprecate-2.c @@ -42,6 +42,3 @@ void f() // { dg-warning "42: the specification of arguments to 'uses_allocators' where each item is of the form 'allocator\\(traits\\)' is deprecated since OpenMP 5.2 \\\[-Wdeprecated-openmp\\\]" "" { target *-*-* } 34 } } - -// { dg-excess-errors "sorry, unimplemented: 'uses_allocators' clause" } -// { dg-excess-errors "sorry, unimplemented: 'uses_allocators' clause" } diff --git a/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 b/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 index 8bc6b768778a..0463f0e0af98 100644 --- a/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 @@ -24,6 +24,10 @@ module omp_lib_kinds parameter :: omp_pteam_mem_alloc = 7 integer (kind=omp_allocator_handle_kind), & parameter :: omp_thread_mem_alloc = 8 + + integer, parameter :: omp_memspace_handle_kind = c_intptr_t + integer (omp_memspace_handle_kind), & + parameter :: omp_default_mem_space = 0 end module subroutine bar (a, b, c) @@ -80,7 +84,8 @@ subroutine foo(x, y) !$omp target teams distribute parallel do private (x) firstprivate (y) & !$omp allocate ((omp_default_mem_alloc + 0):z) allocate & - !$omp (omp_default_mem_alloc: x, y) allocate (h: r) lastprivate (z) reduction(+:r) + !$omp (omp_default_mem_alloc: x, y) allocate (h: r) lastprivate (z) reduction(+:r) & + !$omp uses_allocators(memspace(omp_default_mem_space) : h) do i = 1, 10 call bar (0, x, z); call bar2 (1, y, r); diff --git a/gcc/testsuite/gfortran.dg/gomp/scope-6.f90 b/gcc/testsuite/gfortran.dg/gomp/scope-6.f90 index 4c4f5e034f74..39a65904c337 100644 --- a/gcc/testsuite/gfortran.dg/gomp/scope-6.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/scope-6.f90 @@ -20,4 +20,4 @@ contains end end -! { dg-final { scan-tree-dump "omp scope private\\(a\\) firstprivate\\(b\\) reduction\\(\\+:c\\) allocate\\(allocator\\(D\\.\[0-9\]+\\):a\\) allocate\\(allocator\\(D\\.\[0-9\]+\\):b\\) allocate\\(allocator\\(D\\.\[0-9\]+\\):c\\)" "original" } } +! { dg-final { scan-tree-dump "omp scope private\\(a\\) firstprivate\\(b\\) reduction\\(\\+:c\\) allocate\\(allocator\\(h\\):a\\) allocate\\(allocator\\(h\\):b\\) allocate\\(allocator\\(h\\):c\\)" "original" } } diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index 9d3cb4d81bd3..b470a9dee7b5 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -1155,6 +1155,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_POP_MAPPER_NAME: pp_string (pp, "pop_mapper"); break; + case GOMP_MAP_USES_ALLOCATORS: + pp_string (pp, "uses_allocators"); + break; default: gcc_unreachable (); } diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 32bdedc9c4a8..8662872de7ca 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -195,6 +195,11 @@ enum gomp_map_kind GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION = (GOMP_MAP_DEEP_COPY | 2), + /* Used to pass a descriptor for an OpenMP uses_allocators clause, + also serves as an effective 'firstprivate' clause for the allocator + variable. */ + GOMP_MAP_USES_ALLOCATORS = (GOMP_MAP_DEEP_COPY | 3), + /* Internal to GCC, not used in libgomp. */ /* Do not map, but pointer assign a pointer instead. */ GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1), diff --git a/libgomp/allocator.c b/libgomp/allocator.c index d2e88d97846c..5576d92e21d2 100644 --- a/libgomp/allocator.c +++ b/libgomp/allocator.c @@ -247,14 +247,24 @@ struct omp_allocator_data unsigned int fallback : 8; unsigned int pinned : 1; unsigned int partition : 7; -#if defined(LIBGOMP_USE_MEMKIND) || defined(LIBGOMP_USE_LIBNUMA) + /* To unify the format of this type across host/accelerator, enable + this field unconditionally when offload is enabled. */ + #if defined(LIBGOMP_USE_MEMKIND) || defined(LIBGOMP_USE_LIBNUMA) || \ + defined(LIBGOMP_OFFLOAD_HOST) || defined(LIBGOMP_OFFLOADED_ONLY) unsigned int memkind : 8; #endif + /* Note: we now require __sync builtins for offload host/accelerator, + checked during configuration. This lock should never be enabled + for offload configs. */ #ifndef HAVE_SYNC_BUILTINS gomp_mutex_t lock; #endif }; +/* Size of allocator data, exported within libgomp. */ +const size_t gomp_omp_allocator_data_size attribute_hidden + = sizeof (struct omp_allocator_data); + struct omp_mem_header { void *ptr; @@ -404,38 +414,44 @@ gomp_get_memkind (void) } #endif -omp_allocator_handle_t -omp_init_allocator (omp_memspace_handle_t memspace, int ntraits, - const omp_alloctrait_t traits[]) +static omp_allocator_handle_t +gomp_init_allocator_data (struct omp_allocator_data *data, + omp_memspace_handle_t memspace, int ntraits, + const omp_alloctrait_t traits[]) { - struct omp_allocator_data data - = { memspace, 1, ~(uintptr_t) 0, 0, 0, omp_atv_contended, omp_atv_all, - omp_atv_default_mem_fb, omp_atv_false, omp_atv_environment, -#if defined(LIBGOMP_USE_MEMKIND) || defined(LIBGOMP_USE_LIBNUMA) - GOMP_MEMKIND_NONE + data->memspace = memspace; + data->alignment = 1; + data->pool_size = ~(uintptr_t) 0; + data->used_pool_size = 0; + data->fb_data = 0; + data->sync_hint = omp_atv_contended; + data->access = omp_atv_all; + data->fallback = omp_atv_default_mem_fb; + data->pinned = omp_atv_false; + data->partition = omp_atv_environment; +#if defined(LIBGOMP_USE_MEMKIND) || defined(LIBGOMP_USE_LIBNUMA) || \ + defined(LIBGOMP_OFFLOAD_HOST) || defined(LIBGOMP_OFFLOADED_ONLY) + data->memkind = GOMP_MEMKIND_NONE; #endif - }; - struct omp_allocator_data *ret; - int i; if (memspace > omp_max_predefined_mem_space && (memspace < ompx_gnu_min_predefined_mem_space || memspace > ompx_gnu_max_predefined_mem_space)) return omp_null_allocator; - for (i = 0; i < ntraits; i++) + for (int i = 0; i < ntraits; i++) switch (traits[i].key) { case omp_atk_sync_hint: switch (traits[i].value) { case omp_atv_default: - data.sync_hint = omp_atv_contended; + data->sync_hint = omp_atv_contended; break; case omp_atv_contended: case omp_atv_uncontended: case omp_atv_serialized: case omp_atv_private: - data.sync_hint = traits[i].value; + data->sync_hint = traits[i].value; break; default: return omp_null_allocator; @@ -444,25 +460,25 @@ omp_init_allocator (omp_memspace_handle_t memspace, int ntraits, case omp_atk_alignment: if (traits[i].value == omp_atv_default) { - data.alignment = 1; + data->alignment = 1; break; } if ((traits[i].value & (traits[i].value - 1)) != 0 || !traits[i].value) return omp_null_allocator; - data.alignment = traits[i].value; + data->alignment = traits[i].value; break; case omp_atk_access: switch (traits[i].value) { case omp_atv_default: - data.access = omp_atv_all; + data->access = omp_atv_all; break; case omp_atv_all: case omp_atv_cgroup: case omp_atv_pteam: case omp_atv_thread: - data.access = traits[i].value; + data->access = traits[i].value; break; default: return omp_null_allocator; @@ -470,38 +486,38 @@ omp_init_allocator (omp_memspace_handle_t memspace, int ntraits, break; case omp_atk_pool_size: if (traits[i].value == omp_atv_default) - data.pool_size = ~(uintptr_t) 0; + data->pool_size = ~(uintptr_t) 0; else - data.pool_size = traits[i].value; + data->pool_size = traits[i].value; break; case omp_atk_fallback: switch (traits[i].value) { case omp_atv_default: - data.fallback = omp_atv_default_mem_fb; + data->fallback = omp_atv_default_mem_fb; break; case omp_atv_default_mem_fb: case omp_atv_null_fb: case omp_atv_abort_fb: case omp_atv_allocator_fb: - data.fallback = traits[i].value; + data->fallback = traits[i].value; break; default: return omp_null_allocator; } break; case omp_atk_fb_data: - data.fb_data = traits[i].value; + data->fb_data = traits[i].value; break; case omp_atk_pinned: switch (traits[i].value) { case omp_atv_default: case omp_atv_false: - data.pinned = omp_atv_false; + data->pinned = omp_atv_false; break; case omp_atv_true: - data.pinned = omp_atv_true; + data->pinned = omp_atv_true; break; default: return omp_null_allocator; @@ -511,13 +527,13 @@ omp_init_allocator (omp_memspace_handle_t memspace, int ntraits, switch (traits[i].value) { case omp_atv_default: - data.partition = omp_atv_environment; + data->partition = omp_atv_environment; break; case omp_atv_environment: case omp_atv_nearest: case omp_atv_blocked: case omp_atv_interleaved: - data.partition = traits[i].value; + data->partition = traits[i].value; break; default: return omp_null_allocator; @@ -527,8 +543,8 @@ omp_init_allocator (omp_memspace_handle_t memspace, int ntraits, return omp_null_allocator; } - if (data.alignment < sizeof (void *)) - data.alignment = sizeof (void *); + if (data->alignment < sizeof (void *)) + data->alignment = sizeof (void *); switch (memspace) { @@ -536,57 +552,75 @@ omp_init_allocator (omp_memspace_handle_t memspace, int ntraits, case omp_high_bw_mem_space: struct gomp_memkind_data *memkind_data; memkind_data = gomp_get_memkind (); - if (data.partition == omp_atv_interleaved + if (data->partition == omp_atv_interleaved && memkind_data->kinds[GOMP_MEMKIND_HBW_INTERLEAVE]) { - data.memkind = GOMP_MEMKIND_HBW_INTERLEAVE; + data->memkind = GOMP_MEMKIND_HBW_INTERLEAVE; break; } else if (memkind_data->kinds[GOMP_MEMKIND_HBW_PREFERRED]) { - data.memkind = GOMP_MEMKIND_HBW_PREFERRED; + data->memkind = GOMP_MEMKIND_HBW_PREFERRED; break; } break; case omp_large_cap_mem_space: memkind_data = gomp_get_memkind (); if (memkind_data->kinds[GOMP_MEMKIND_DAX_KMEM_ALL]) - data.memkind = GOMP_MEMKIND_DAX_KMEM_ALL; + data->memkind = GOMP_MEMKIND_DAX_KMEM_ALL; else if (memkind_data->kinds[GOMP_MEMKIND_DAX_KMEM]) - data.memkind = GOMP_MEMKIND_DAX_KMEM; + data->memkind = GOMP_MEMKIND_DAX_KMEM; break; #endif default: #ifdef LIBGOMP_USE_MEMKIND - if (data.partition == omp_atv_interleaved) + if (data->partition == omp_atv_interleaved) { memkind_data = gomp_get_memkind (); if (memkind_data->kinds[GOMP_MEMKIND_INTERLEAVE]) - data.memkind = GOMP_MEMKIND_INTERLEAVE; + data->memkind = GOMP_MEMKIND_INTERLEAVE; } #endif break; } #ifdef LIBGOMP_USE_LIBNUMA - if (data.memkind == GOMP_MEMKIND_NONE && data.partition == omp_atv_nearest) + if (data->memkind == GOMP_MEMKIND_NONE && data->partition == omp_atv_nearest) { libnuma_data = gomp_get_libnuma (); if (libnuma_data->numa_alloc_local != NULL) - data.memkind = GOMP_MEMKIND_LIBNUMA; + data->memkind = GOMP_MEMKIND_LIBNUMA; } #endif - /* Reject unsupported memory spaces. */ - if (!MEMSPACE_VALIDATE (data.memspace, data.access, data.pinned)) - return omp_null_allocator; + return (omp_allocator_handle_t) data; +} - ret = gomp_malloc (sizeof (struct omp_allocator_data)); - *ret = data; +omp_allocator_handle_t +omp_init_allocator (omp_memspace_handle_t memspace, int ntraits, + const omp_alloctrait_t traits[]) +{ + struct omp_allocator_data *data + = gomp_malloc (sizeof (struct omp_allocator_data)); + + omp_allocator_handle_t ret + = gomp_init_allocator_data (data, memspace, ntraits, traits); + + if (ret == omp_null_allocator + /* Reject unsupported memory spaces. */ + || !MEMSPACE_VALIDATE (data->memspace, data->access, data->pinned)) + { + ret = omp_null_allocator; + free (data); + } + else + { #ifndef HAVE_SYNC_BUILTINS - gomp_mutex_init (&ret->lock); + gomp_mutex_init (data->lock); #endif - return (omp_allocator_handle_t) ret; + } + + return ret; } void @@ -1500,3 +1534,61 @@ fail:; } return NULL; } + +#if !defined(LIBGOMP_OFFLOADED_ONLY) + +/* Called from gomp_maps_vars, used to implement the effective + omp_init_allocator() call in an uses_allocators clause. + + The memspace_validate hook is also used here to check things for + the offload accelerator on the host side. +*/ +attribute_hidden uintptr_t +gomp_map_omp_init_allocator (struct gomp_device_descr *devicep, + struct goacc_asyncqueue *aq, + struct gomp_coalesce_buf *cbuf, + void *allocator_data_devaddr, void *descr_ptr) +{ + uintptr_t *descr = (uintptr_t *) descr_ptr; + + omp_memspace_handle_t memspace = (omp_memspace_handle_t) descr[0]; + int ntraits = (int) descr[1]; + omp_alloctrait_t *traits = (omp_alloctrait_t *) descr[2]; + + struct omp_allocator_data data; + omp_allocator_handle_t ret = gomp_init_allocator_data (&data, memspace, + ntraits, traits); + if (ret != omp_null_allocator) + { + if (devicep) + { + /* Do memspace validation for the offload target, using the + offload plugin. */ + if (devicep->memspace_validate_func + && !devicep->memspace_validate_func (data.memspace, data.access)) + return (uintptr_t) omp_null_allocator; + + /* Copy to device. This is now the device-side omp_allocator_data. */ + gomp_copy_host2dev (devicep, aq, allocator_data_devaddr, &data, + sizeof (struct omp_allocator_data), true, cbuf); + } + else + { + /* Used for host fallback case. */ + + if (!MEMSPACE_VALIDATE (data.memspace, data.access, data.pinned)) + /* Reject unsupported memory spaces. */ + return omp_null_allocator; + + memcpy (allocator_data_devaddr, &data, + sizeof (struct omp_allocator_data)); +#ifndef HAVE_SYNC_BUILTINS + gomp_mutex_init + (((struct omp_allocator_data *) allocator_data_devaddr)->lock); +#endif + } + ret = (omp_allocator_handle_t) allocator_data_devaddr; + } + return (uintptr_t) ret; +} +#endif diff --git a/libgomp/config/gcn/allocator.c b/libgomp/config/gcn/allocator.c index b3d7bcef7d9d..f09d19cc755e 100644 --- a/libgomp/config/gcn/allocator.c +++ b/libgomp/config/gcn/allocator.c @@ -110,6 +110,10 @@ gcn_memspace_realloc (omp_memspace_handle_t memspace, void *addr, return realloc (addr, size); } +/* The MEMSPACE_VALIDATE implementation should be kept in sync with + the GOMP_OFFLOAD_memspace_validate plugin hook function in + plugin/plugin-gcn.c. */ + static inline int gcn_memspace_validate (omp_memspace_handle_t memspace, unsigned access) { diff --git a/libgomp/config/nvptx/allocator.c b/libgomp/config/nvptx/allocator.c index ebbe98169098..577ebb7d1a58 100644 --- a/libgomp/config/nvptx/allocator.c +++ b/libgomp/config/nvptx/allocator.c @@ -54,6 +54,9 @@ asm (".extern .shared .u8 __nvptx_lowlat_pool[];\n"); static void * nvptx_memspace_alloc (omp_memspace_handle_t memspace, size_t size) { +#if __PTX_ISA_VERSION_MAJOR__ > 4 \ + || (__PTX_ISA_VERSION_MAJOR__ == 4 && __PTX_ISA_VERSION_MINOR >= 1) + /* Low-latency memory is not available before PTX 4.1. */ if (memspace == omp_low_lat_mem_space) { char *shared_pool; @@ -65,12 +68,16 @@ nvptx_memspace_alloc (omp_memspace_handle_t memspace, size_t size) /* No non-standard memspaces are implemented for device-side nvptx. */ return NULL; else +#endif return malloc (size); } static void * nvptx_memspace_calloc (omp_memspace_handle_t memspace, size_t size) { +#if __PTX_ISA_VERSION_MAJOR__ > 4 \ + || (__PTX_ISA_VERSION_MAJOR__ == 4 && __PTX_ISA_VERSION_MINOR >= 1) + /* Low-latency memory is not available before PTX 4.1. */ if (memspace == omp_low_lat_mem_space) { char *shared_pool; @@ -82,12 +89,16 @@ nvptx_memspace_calloc (omp_memspace_handle_t memspace, size_t size) /* No non-standard memspaces are implemented for device-side nvptx. */ return NULL; else +#endif return calloc (1, size); } static void nvptx_memspace_free (omp_memspace_handle_t memspace, void *addr, size_t size) { +#if __PTX_ISA_VERSION_MAJOR__ > 4 \ + || (__PTX_ISA_VERSION_MAJOR__ == 4 && __PTX_ISA_VERSION_MINOR >= 1) + /* Low-latency memory is not available before PTX 4.1. */ if (memspace == omp_low_lat_mem_space) { char *shared_pool; @@ -96,6 +107,7 @@ nvptx_memspace_free (omp_memspace_handle_t memspace, void *addr, size_t size) __nvptx_lowlat_free (shared_pool, addr, size); } else +#endif free (addr); } @@ -103,6 +115,9 @@ static void * nvptx_memspace_realloc (omp_memspace_handle_t memspace, void *addr, size_t oldsize, size_t size) { +#if __PTX_ISA_VERSION_MAJOR__ > 4 \ + || (__PTX_ISA_VERSION_MAJOR__ == 4 && __PTX_ISA_VERSION_MINOR >= 1) + /* Low-latency memory is not available before PTX 4.1. */ if (memspace == omp_low_lat_mem_space) { char *shared_pool; @@ -110,13 +125,21 @@ nvptx_memspace_realloc (omp_memspace_handle_t memspace, void *addr, return __nvptx_lowlat_realloc (shared_pool, addr, oldsize, size); } - else if (memspace > GOMP_OMP_PREDEF_MEMSPACE_MAX) - /* No non-standard memspaces are implemented for device-side nvptx. */ - return NULL; else - return realloc (addr, size); +#endif + { + if (memspace > GOMP_OMP_PREDEF_MEMSPACE_MAX) + /* No non-standard memspaces are implemented for device-side nvptx. */ + return NULL; + else + return realloc (addr, size); + } } +/* The MEMSPACE_VALIDATE implementation should be kept in sync with + the GOMP_OFFLOAD_memspace_validate plugin hook function in + plugin/plugin-nvptx.c. */ + static inline int nvptx_memspace_validate (omp_memspace_handle_t memspace, unsigned access) { diff --git a/libgomp/configure b/libgomp/configure index f7b510a7200b..7877089755a5 100755 --- a/libgomp/configure +++ b/libgomp/configure @@ -15782,6 +15782,12 @@ fi +if test x"$offload_plugins" != x; then + +$as_echo "#define LIBGOMP_OFFLOAD_HOST 1" >>confdefs.h + +fi + # Check for functions needed. for ac_func in getloadavg clock_gettime strtoull do : @@ -17120,6 +17126,19 @@ $as_echo "#define HAVE_SYNC_BUILTINS 1" >>confdefs.h fi +# If offloading is enabled, both offload hosts and accelerators are required to support +# __sync_* builtins, to simplify architectural assumptions between both sides +if test x$libgomp_cv_have_sync_builtins = xno; then + # We require accelerator targets to support __sync_* builtins. + if test x$libgomp_offloaded_only = xyes; then + as_fn_error $? "accelerator targets require __sync_val_compare_and_swap to build libgomp." "$LINENO" 5 + fi + # Same for offload hosts. + if test x"$offload_plugins" = x; then + as_fn_error $? "offload hosts require __sync_val_compare_and_swap to build libgomp." "$LINENO" 5 + fi +fi + XCFLAGS="$XCFLAGS$XPCFLAGS" # Add CET specific flags if CET is enabled diff --git a/libgomp/configure.ac b/libgomp/configure.ac index 8245c737d42a..917079bc95bc 100644 --- a/libgomp/configure.ac +++ b/libgomp/configure.ac @@ -233,6 +233,11 @@ AC_CHECK_SIZEOF([void *]) m4_include([plugin/configfrag.ac]) +if test x"$offload_plugins" != x; then + AC_DEFINE(LIBGOMP_OFFLOAD_HOST, 1, + [Define to 1 if building libgomp for a offload host]) +fi + # Check for functions needed. AC_CHECK_FUNCS(getloadavg clock_gettime strtoull) AC_CHECK_FUNCS(aligned_alloc posix_memalign memalign _aligned_malloc) @@ -365,6 +370,19 @@ CFLAGS="$save_CFLAGS $XCFLAGS" # had a chance to set XCFLAGS. LIBGOMP_CHECK_SYNC_BUILTINS +# If offloading is enabled, both offload hosts and accelerators are required to support +# __sync_* builtins, to simplify architectural assumptions between both sides +if test x$libgomp_cv_have_sync_builtins = xno; then + # We require accelerator targets to support __sync_* builtins. + if test x$libgomp_offloaded_only = xyes; then + AC_MSG_ERROR([accelerator targets require __sync_val_compare_and_swap to build libgomp.]) + fi + # Same for offload hosts. + if test x"$offload_plugins" = x; then + AC_MSG_ERROR([offload hosts require __sync_val_compare_and_swap to build libgomp.]) + fi +fi + XCFLAGS="$XCFLAGS$XPCFLAGS" # Add CET specific flags if CET is enabled diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h index 60ac7201597a..db7c03fb350b 100644 --- a/libgomp/libgomp-plugin.h +++ b/libgomp/libgomp-plugin.h @@ -191,6 +191,7 @@ extern int GOMP_OFFLOAD_memcpy3d (int, int, size_t, size_t, size_t, void *, const void *, size_t, size_t, size_t, size_t, size_t); extern bool GOMP_OFFLOAD_memset (int, void *, int, size_t); +extern int GOMP_OFFLOAD_memspace_validate (omp_memspace_handle_t, unsigned int); extern bool GOMP_OFFLOAD_can_run (void *); /* An opaque type, encapsulating the state required to launch a single diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 9510b148ed8c..3c8a2a1115e6 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -616,6 +616,7 @@ extern size_t gomp_affinity_format_len; extern int gomp_get_current_numa_node (); extern int gomp_get_numa_distance (int, int); extern uintptr_t gomp_def_allocator; +extern const size_t gomp_omp_allocator_data_size; extern const struct gomp_default_icv gomp_default_icv_values; extern struct gomp_icv_list *gomp_initial_icv_list; extern struct gomp_offload_icv_list *gomp_offload_icv_list; @@ -1472,6 +1473,7 @@ struct gomp_device_descr __typeof (GOMP_OFFLOAD_memcpy2d) *memcpy2d_func; __typeof (GOMP_OFFLOAD_memcpy3d) *memcpy3d_func; __typeof (GOMP_OFFLOAD_memset) *memset_func; + __typeof (GOMP_OFFLOAD_memspace_validate) *memspace_validate_func; struct { __typeof (GOMP_OFFLOAD_session_start) *start_func; __typeof (GOMP_OFFLOAD_session_allocate_target_var_table) *alloc_tvt_func; @@ -1561,6 +1563,13 @@ extern bool gomp_remove_var (struct gomp_device_descr *, splay_tree_key); extern void gomp_remove_var_async (struct gomp_device_descr *, splay_tree_key, struct goacc_asyncqueue *); +/* allocator.c */ + +extern uintptr_t gomp_map_omp_init_allocator (struct gomp_device_descr *, + struct goacc_asyncqueue *, + struct gomp_coalesce_buf *, + void *, void *); + /* work.c */ extern void gomp_init_work_share (struct gomp_work_share *, size_t, unsigned); diff --git a/libgomp/oacc-host.c b/libgomp/oacc-host.c index 2cc7e6464eae..b0546c87405f 100644 --- a/libgomp/oacc-host.c +++ b/libgomp/oacc-host.c @@ -320,6 +320,7 @@ static struct gomp_device_descr host_dispatch = .host2dev_func = host_host2dev, .memcpy2d_func = NULL, .memcpy3d_func = NULL, + .memspace_validate_func = NULL, .run_func = host_run, .session = { diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c index 81426168da95..cc47d205e1e6 100644 --- a/libgomp/plugin/plugin-gcn.c +++ b/libgomp/plugin/plugin-gcn.c @@ -4896,6 +4896,22 @@ unlock: return retval; } +/* This plugin hook function should be kept in sync with gcn_memspace_validate + in config/gcn/allocator.c. */ + +int +GOMP_OFFLOAD_memspace_validate (omp_memspace_handle_t memspace, unsigned access) +{ + /* Disallow use of low-latency memory when it must be accessible by + all threads. */ + return (memspace != omp_low_lat_mem_space + || access != omp_atv_all); + + /* Otherwise, standard memspaces are accepted, even when we don't have + anything special to do with them, and non-standard memspaces are assumed + to need explicit support. */ + return (memspace <= GOMP_OMP_PREDEF_MEMSPACE_MAX); +} static bool init_hip_runtime_functions (void) diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c index ddb699a302f8..c85ff89002e1 100644 --- a/libgomp/plugin/plugin-nvptx.c +++ b/libgomp/plugin/plugin-nvptx.c @@ -2583,6 +2583,24 @@ GOMP_OFFLOAD_memset (int ord, void *ptr, int val, size_t count) return true; } +/* This plugin hook function should be kept in sync with nvptx_memspace_validate + in config/nvptx/allocator.c. */ + +int +GOMP_OFFLOAD_memspace_validate (omp_memspace_handle_t memspace, unsigned access) +{ + /* Disallow use of low-latency memory when it must be accessible by + all threads. */ + if (memspace == omp_low_lat_mem_space + && access == omp_atv_all) + return false; + + /* Otherwise, standard memspaces are accepted, even when we don't have + anything special to do with them, and non-standard memspaces are assumed + to need explicit support. */ + return (memspace <= GOMP_OMP_PREDEF_MEMSPACE_MAX); +} + bool GOMP_OFFLOAD_openacc_async_host2dev (int ord, void *dst, const void *src, size_t n, struct goacc_asyncqueue *aq) diff --git a/libgomp/target.c b/libgomp/target.c index 78f390d593f6..e49eacc00713 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1322,6 +1322,23 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, { tgt->list[i].key = NULL; tgt->list[i].offset = OFFSET_INLINED; + has_firstprivate = true; + continue; + } + else if ((kind & typemask) == GOMP_MAP_USES_ALLOCATORS) + { + tgt->list[i].key = NULL; + tgt->list[i].offset = OFFSET_INLINED; + + size_t align = (size_t) 1 << (kind >> rshift); + if (tgt_align < align) + tgt_align = align; + tgt_size = (tgt_size + align - 1) & ~(align - 1); + + /* Allocate space for omp_allocator_data. */ + tgt_size += gomp_omp_allocator_data_size; + + has_firstprivate = true; continue; } else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR @@ -1753,6 +1770,21 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, tgt->list[i].offset = OFFSET_INLINED; } continue; + + case GOMP_MAP_USES_ALLOCATORS: + align = (size_t) 1 << (kind >> rshift); + tgt_size = (tgt_size + align - 1) & ~(align - 1); + tgt->list[i].offset = tgt_size; + + void *descr_ptr = hostaddrs[i]; + void *devaddr = (void *) (tgt->tgt_start + tgt_size); + uintptr_t dev_allocator + = gomp_map_omp_init_allocator (devicep, aq, cbufp, devaddr, + descr_ptr); + hostaddrs[i] = (void *) dev_allocator; + tgt_size += gomp_omp_allocator_data_size; + continue; + case GOMP_MAP_STRUCT_UNORD: if (sizes[i] > 1) { @@ -3204,6 +3236,14 @@ calculate_firstprivate_requirements (size_t mapnum, size_t *sizes, *tgt_size = (*tgt_size + align - 1) & ~(align - 1); *tgt_size += sizes[i]; } + else if ((kinds[i] & 0xff) == GOMP_MAP_USES_ALLOCATORS) + { + size_t align = (size_t) 1 << (kinds[i] >> 8); + if (*tgt_align < align) + *tgt_align = align; + *tgt_size = (*tgt_size + align - 1) & ~(align - 1); + *tgt_size += gomp_omp_allocator_data_size; + } } /* Copy data shared as GOMP_MAP_FIRSTPRIVATE to DST. */ @@ -3226,12 +3266,23 @@ copy_firstprivate_data (char *tgt, size_t mapnum, void **hostaddrs, memcpy (tgt + tgt_size, hostaddrs[i], sizes[i]); hostaddrs[i] = tgt + tgt_size; tgt_size = tgt_size + sizes[i]; - if (i + 1 < mapnum && (kinds[i+1] & 0xff) == GOMP_MAP_ATTACH) + if (i + 1 < mapnum && (kinds[i + 1] & 0xff) == GOMP_MAP_ATTACH) { - *(*(uintptr_t**) hostaddrs[i+1] + sizes[i+1]) = (uintptr_t) hostaddrs[i]; + *(*(uintptr_t**) hostaddrs[i + 1] + sizes[i + 1]) + = (uintptr_t) hostaddrs[i]; ++i; } } + else if ((kinds[i] & 0xff) == GOMP_MAP_USES_ALLOCATORS) + { + assert (hostaddrs[i] != NULL); + size_t align = (size_t) 1 << (kinds[i] >> 8); + tgt_size = (tgt_size + align - 1) & ~(align - 1); + uintptr_t allocator = gomp_map_omp_init_allocator (NULL, NULL, NULL, + tgt + tgt_size, + hostaddrs[i]); + hostaddrs[i] = (void *) allocator; + } } /* Helper function of GOMP_target{,_ext} routines. */ @@ -6261,6 +6312,7 @@ gomp_load_plugin_for_device (struct gomp_device_descr *device, DLSYM (host2dev); DLSYM_OPT (memcpy2d, memcpy2d); DLSYM_OPT (memcpy3d, memcpy3d); + DLSYM_OPT (memspace_validate, memspace_validate); if (DLSYM_OPT (interop, interop)) { DLSYM (get_interop_int); diff --git a/libgomp/testsuite/libgomp.c++/c++.exp b/libgomp/testsuite/libgomp.c++/c++.exp index ed096e17b9c5..5be949bb611d 100644 --- a/libgomp/testsuite/libgomp.c++/c++.exp +++ b/libgomp/testsuite/libgomp.c++/c++.exp @@ -1,6 +1,15 @@ load_lib libgomp-dg.exp load_gcc_lib gcc-dg.exp +proc check_effective_target_c { } { + return 0 +} + +proc check_effective_target_c++ { } { + return 1 +} + + if { $blddir != "" } { set libstdc++_library_path "../libstdc++-v3/src/.libs" set shlib_ext [get_shlib_extension] diff --git a/gcc/testsuite/g++.dg/gomp/uses_allocators-1.C b/libgomp/testsuite/libgomp.c++/uses_allocators-1.C similarity index 89% rename from gcc/testsuite/g++.dg/gomp/uses_allocators-1.C rename to libgomp/testsuite/libgomp.c++/uses_allocators-1.C index 79e51638534c..79b5cb0883f0 100644 --- a/gcc/testsuite/g++.dg/gomp/uses_allocators-1.C +++ b/libgomp/testsuite/libgomp.c++/uses_allocators-1.C @@ -1,8 +1,8 @@ // { dg-do compile } /* { dg-additional-options "-Wno-deprecated-openmp" } */ -//#include <omp.h> - +#include <omp.h> +/* typedef __UINTPTR_TYPE__ omp_uintptr_t; #if __cplusplus >= 201103L @@ -43,7 +43,7 @@ typedef struct omp_alloctrait_t // omp_alloctrait_key_t key; // omp_uintptr_t value; } omp_alloctrait_t; - +*/ template<typename TH, TH alloc> void f() @@ -99,7 +99,3 @@ void use2() g2<omp_allocator_handle_t, const omp_alloctrait_t[]>(my); // OK } - -// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 51 } -// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 58 } -// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 89 } diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-1.c similarity index 64% rename from gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c rename to libgomp/testsuite/libgomp.c-c++-common/uses_allocators-1.c index df82cbbcba9d..2029ee5a9a3a 100644 --- a/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c +++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-1.c @@ -1,31 +1,16 @@ -typedef enum omp_allocator_handle_t -#if __cplusplus >= 201103L -: __UINTPTR_TYPE__ -#endif -{ - omp_default_mem_alloc = 1, - omp_low_lat_mem_alloc = 5, - __omp_allocator_handle_t_max__ = __UINTPTR_MAX__ -} omp_allocator_handle_t; - -typedef struct omp_alloctrait_t -{ - int key; - int value; -} omp_alloctrait_t; - -extern void *omp_alloc (__SIZE_TYPE__, omp_allocator_handle_t); +/* { dg-do compile } */ +#include <omp.h> void f (omp_allocator_handle_t my_alloc) { #pragma omp target { - int a; /* { dg-error "'my_alloc' in 'allocator' clause inside a target region must be specified in an 'uses_allocators' clause on the 'target' directive" "not yet implemented" { xfail *-*-* } } */ + int a; /* { dg-error "'my_alloc' in 'allocator' clause inside a target region must be specified in an 'uses_allocators' clause on the 'target' directive" "" { target c } } */ #pragma omp allocate(a) allocator(my_alloc) /* { dg-message "sorry, unimplemented: '#pragma omp allocate' not yet supported" "" { target c++ } } */ a = 5; void *prt = omp_alloc(32, my_alloc); - #pragma omp parallel allocate(allocator(my_alloc) : a) firstprivate(a) /* { dg-error "allocator 'my_alloc' in 'allocate' clause inside a target region must be specified in an 'uses_allocators' clause on the 'target' directive" "not yet implemented" { xfail *-*-* } } */ + #pragma omp parallel allocate(allocator(my_alloc) : a) firstprivate(a) /* { dg-error "allocator 'my_alloc' in 'allocate' clause inside a target region must be specified in an 'uses_allocators' clause on the 'target' directive" } */ a = 7; } } @@ -44,5 +29,3 @@ g (omp_allocator_handle_t my_alloc) a = 7; } } - -// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 37 } diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-2.c similarity index 51% rename from gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c rename to libgomp/testsuite/libgomp.c-c++-common/uses_allocators-2.c index c766ea3559d0..2babb11018e2 100644 --- a/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c +++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-2.c @@ -1,18 +1,5 @@ -typedef enum omp_allocator_handle_t -#if __cplusplus >= 201103L -: __UINTPTR_TYPE__ -#endif -{ - omp_default_mem_alloc = 1, - omp_low_lat_mem_alloc = 5, - __omp_allocator_handle_t_max__ = __UINTPTR_MAX__ -} omp_allocator_handle_t; - -typedef struct omp_alloctrait_t -{ - int key; - int value; -} omp_alloctrait_t; +/* { dg-do compile } */ +#include <omp.h> void f () @@ -31,5 +18,3 @@ g () #pragma omp target uses_allocators(traits(trait) : my_alloc) ; } - -// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 31 } diff --git a/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-3.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-3.c new file mode 100644 index 000000000000..a371d7fe8238 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-3.c @@ -0,0 +1,50 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */ + +#include <omp.h> + +int n = 2; + +int main (void) +{ + omp_allocator_handle_t foo, bar; + const omp_alloctrait_t foo_traits[] = { { omp_atk_pinned, omp_atv_true }, + { omp_atk_partition, omp_atv_nearest } }; + + #pragma omp target + ; + #pragma omp target uses_allocators (bar) + ; + #pragma omp target uses_allocators (foo (foo_traits)) /* { dg-warning "the specification of arguments to 'uses_allocators' where each item is of the form 'allocator\\(traits\\)' is deprecated since OpenMP 5.2 \\\[-Wdeprecated-openmp\\\]" } */ + ; + #pragma omp target uses_allocators (foo (foo_traits), bar (foo_traits)) /* { dg-warning "the specification of arguments to 'uses_allocators' where each item is of the form 'allocator\\(traits\\)' is deprecated since OpenMP 5.2 \\\[-Wdeprecated-openmp\\\]" } */ + ; + #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo) + ; + #pragma omp target uses_allocators (traits(foo_traits) : bar) + ; + #pragma omp target parallel uses_allocators (memspace(omp_high_bw_mem_space), traits(foo_traits) : bar) + ; + #pragma omp target parallel uses_allocators (traits(foo_traits), memspace(omp_high_bw_mem_space) : bar) uses_allocators(foo) + { + void *p = omp_alloc ((unsigned long) 32, bar); + omp_free (p, bar); + } + return 0; +} + +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(\\), traits\\(\\) : bar\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(\\), traits\\(foo_traits\\) : foo\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(\\), traits\\(foo_traits\\) : bar\\) uses_allocators\\(memspace\\(\\), traits\\(foo_traits\\) : foo\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(.+\\), traits\\(\\) : foo\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(\\), traits\\(foo_traits\\) : bar\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(.+\\), traits\\(foo_traits\\) : bar\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(.+\\), traits\\(foo_traits\\) : bar\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : foo\\)" "original" } } */ + +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : bar\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(\\), traits\\(foo_traits\\) : foo\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(\\), traits\\(foo_traits\\) : bar\\) uses_allocators\\(memspace\\(\\), traits\\(foo_traits\\) : foo\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(.+\\), traits\\(\\) : foo\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(\\), traits\\(foo_traits\\) : bar\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(.+\\), traits\\(foo_traits\\) : bar\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(.+\\), traits\\(foo_traits\\) : bar\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : foo\\)" "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-4.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-4.c similarity index 70% rename from gcc/testsuite/c-c++-common/gomp/uses_allocators-4.c rename to libgomp/testsuite/libgomp.c-c++-common/uses_allocators-4.c index 5f3650157fc4..389a51cd1c59 100644 --- a/gcc/testsuite/c-c++-common/gomp/uses_allocators-4.c +++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-4.c @@ -1,86 +1,7 @@ /* { dg-do compile } */ /* { dg-additional-options "-Wno-deprecated-openmp" } */ -//#include <omp.h> - -typedef __UINTPTR_TYPE__ omp_uintptr_t; - -#if __cplusplus >= 201103L -# define __GOMP_UINTPTR_T_ENUM : omp_uintptr_t -#else -# define __GOMP_UINTPTR_T_ENUM -#endif - -typedef enum omp_memspace_handle_t __GOMP_UINTPTR_T_ENUM -{ - omp_default_mem_space = 0, - omp_large_cap_mem_space = 1, - omp_const_mem_space = 2, - omp_high_bw_mem_space = 3, - omp_low_lat_mem_space = 4, - ompx_gnu_managed_mem_space = 200, - __omp_memspace_handle_t_max__ = __UINTPTR_MAX__ -} omp_memspace_handle_t; - -typedef enum omp_allocator_handle_t __GOMP_UINTPTR_T_ENUM -{ - omp_null_allocator = 0, - omp_default_mem_alloc = 1, - omp_large_cap_mem_alloc = 2, - omp_const_mem_alloc = 3, - omp_high_bw_mem_alloc = 4, - omp_low_lat_mem_alloc = 5, - omp_cgroup_mem_alloc = 6, - omp_pteam_mem_alloc = 7, - omp_thread_mem_alloc = 8, - ompx_gnu_pinned_mem_alloc = 200, - ompx_gnu_managed_mem_alloc = 201, - __omp_allocator_handle_t_max__ = __UINTPTR_MAX__ -} omp_allocator_handle_t; - -typedef enum omp_alloctrait_key_t -{ - omp_atk_sync_hint = 1, - omp_atk_alignment = 2, - omp_atk_access = 3, - omp_atk_pool_size = 4, - omp_atk_fallback = 5, - omp_atk_fb_data = 6, - omp_atk_pinned = 7, - omp_atk_partition = 8 -} omp_alloctrait_key_t; - -typedef enum omp_alloctrait_value_t -{ - omp_atv_default = (__UINTPTR_TYPE__) -1, - omp_atv_false = 0, - omp_atv_true = 1, - omp_atv_contended = 3, - omp_atv_uncontended = 4, - omp_atv_serialized = 5, - omp_atv_private = 6, - omp_atv_all = 7, - omp_atv_thread = 8, - omp_atv_pteam = 9, - omp_atv_cgroup = 10, - omp_atv_default_mem_fb = 11, - omp_atv_null_fb = 12, - omp_atv_abort_fb = 13, - omp_atv_allocator_fb = 14, - omp_atv_environment = 15, - omp_atv_nearest = 16, - omp_atv_blocked = 17, - omp_atv_interleaved = 18 -} omp_alloctrait_value_t; - -typedef struct omp_alloctrait_t -{ - omp_alloctrait_key_t key; - omp_uintptr_t value; -} omp_alloctrait_t; - -omp_alloctrait_key_t k; -omp_alloctrait_value_t v; +#include <omp.h> int f (const omp_alloctrait_t arg_traits[], int n) { @@ -135,10 +56,3 @@ int f (const omp_alloctrait_t arg_traits[], int n) /* { dg-error "legacy 'foo\\\(foo_traits\\\)' traits syntax not allowed in 'uses_allocators' clause when using modifiers" "" { target *-*-* } .-2 } */ return 0; } - -// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 103 } -// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target c } 111 } -// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target c } 113 } -// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target c } 117 } -// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target c } 119 } -// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 131 } diff --git a/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-5.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-5.c new file mode 100644 index 000000000000..eeac5f097cb8 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-5.c @@ -0,0 +1,34 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */ + +#include <omp.h> + +int main (void) +{ + omp_allocator_handle_t memspace, traits; + const omp_alloctrait_t mytraits[] = { { omp_atk_pinned, omp_atv_true }, + { omp_atk_partition, omp_atv_nearest } }; + #pragma omp target uses_allocators (memspace) + ; + #pragma omp target uses_allocators (traits) + ; + #pragma omp target uses_allocators (traits, memspace) + ; + #pragma omp target uses_allocators (traits (mytraits)) /* { dg-warning "the specification of arguments to 'uses_allocators' where each item is of the form 'allocator\\(traits\\)' is deprecated since OpenMP 5.2" } */ + ; + #pragma omp target uses_allocators (memspace (mytraits), omp_default_mem_alloc) /* { dg-warning "the specification of arguments to 'uses_allocators' where each item is of the form 'allocator\\(traits\\)' is deprecated since OpenMP 5.2" } */ + ; + return 0; +} + +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(\\), traits\\(\\) : memspace\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(\\), traits\\(\\) : traits\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(\\), traits\\(\\) : memspace\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : traits\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(\\), traits\\(mytraits\\) : traits\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(\\), traits\\(mytraits\\) : memspace\\)" "original" } } */ + +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : memspace\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : traits\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : memspace\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : traits\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(\\), traits\\(mytraits\\) : traits\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(\\), traits\\(mytraits\\) : memspace\\)" "gimple" } } */ diff --git a/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-6.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-6.c new file mode 100644 index 000000000000..23a831e8ba90 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-6.c @@ -0,0 +1,53 @@ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#include <stdint.h> +#include <omp.h> + +int +main () +{ + int x, *xbuf[10]; + omp_allocator_handle_t my_alloc; + const omp_alloctrait_t trait[1]= {{omp_atk_alignment,128}}; + + #pragma omp target uses_allocators(omp_low_lat_mem_alloc) map(tofrom: x, xbuf) defaultmap(none) + #pragma omp parallel allocate(allocator(omp_low_lat_mem_alloc), align(128): x, xbuf) if(0) firstprivate(x, xbuf) + { + if ((uintptr_t) &x % 128 != 0) + __builtin_abort (); + if ((uintptr_t) xbuf % 128 != 0) + __builtin_abort (); + } + + my_alloc = (omp_allocator_handle_t) 0xABCD; + + #pragma omp target uses_allocators(traits(trait): my_alloc) defaultmap(none) map(tofrom: x, xbuf) + #pragma omp parallel allocate(allocator(my_alloc): x, xbuf) if(0) firstprivate(x, xbuf) + { + if ((uintptr_t) &x % 128 != 0) + __builtin_abort (); + if ((uintptr_t) xbuf % 128 != 0) + __builtin_abort (); + } + + if (my_alloc != (omp_allocator_handle_t) 0xABCD) + __builtin_abort (); + + /* The following creates an allocator with empty traits + default mem space. */ + #pragma omp target uses_allocators(my_alloc) map(tofrom: x, xbuf) defaultmap(none) + #pragma omp parallel allocate(allocator(my_alloc), align(128): x, xbuf) if(0) firstprivate(x, xbuf) + { + if ((uintptr_t) &x % 128 != 0) + __builtin_abort (); + if ((uintptr_t) xbuf % 128 != 0) + __builtin_abort (); + } + + if (my_alloc != (omp_allocator_handle_t) 0xABCD) + __builtin_abort (); + + return 0; +} + +/* { dg-final { scan-tree-dump-times "#pragma omp target .* uses_allocators\\(memspace\\(\\), traits\\(trait\\) : my_alloc\\)" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp target .* uses_allocators\\(memspace\\(\\), traits\\(\\) : my_alloc\\)" 1 "gimple" } } */ diff --git a/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-7.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-7.c new file mode 100644 index 000000000000..7603dfcd3226 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-7.c @@ -0,0 +1,35 @@ +/* { dg-do compile } */ +#include <omp.h> + +void f() +{ + omp_allocator_handle_t my; + struct t { + omp_allocator_handle_t a1; + } s; + +const omp_allocator_handle_t my3 = (omp_allocator_handle_t) 300; +omp_allocator_handle_t my4[1]; +const omp_alloctrait_t t[] = {}; + #pragma omp target uses_allocators(my, omp_default_mem_alloc, omp_null_allocator) // OK + ; + #pragma omp target uses_allocators(my) firstprivate(my) // { dg-error "'my' appears more than once in data clauses" } + ; + #pragma omp target private(my) uses_allocators(my) // { dg-error "'my' appears more than once in data clauses" } + ; + #pragma omp target uses_allocators(my3) + ; + #pragma omp target uses_allocators(s.a1) + // { dg-error "expected '\\)' before '\\.' token" "" { target *-*-* } .-1 } + // { dg-error "allocator 's' must be of 'omp_allocator_handle_t' type" "" { target *-*-* } .-2 } + ; + #pragma omp target uses_allocators(my4) + // { dg-error "allocator 'my4' must be of 'omp_allocator_handle_t' type" "" { target *-*-* } .-1 } + ; + #pragma omp target uses_allocators(my4[0]) + // { dg-error "expected '\\)' before '\\\[' token" "" { target *-*-* } .-1 } + // { dg-error "allocator 'my4' must be of 'omp_allocator_handle_t' type" "" { target *-*-* } .-2 } + ; + #pragma omp target uses_allocators(memspace(omp_default_mem_space) : my, my(t)) // { dg-error "legacy 'my\\(t\\)' traits syntax not allowed in 'uses_allocators' clause when using modifiers" } + ; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-8.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-8.c new file mode 100644 index 000000000000..fef13585dbfc --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-8.c @@ -0,0 +1,16 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-original" } + +#include <omp.h> + +void f() +{ + omp_allocator_handle_t my, my2, my3, my4; +const omp_alloctrait_t t[] = {}; +const omp_alloctrait_t t2[] = {}; + #pragma omp target uses_allocators(traits(t), memspace(omp_high_bw_mem_space) : my; omp_default_mem_alloc, omp_null_allocator; my2; traits(t2) : my3; memspace(omp_large_cap_mem_space) : my4) + ; +} + +// { dg-final { scan-tree-dump "#pragma omp target uses_allocators\\(memspace\\(1\\), traits\\(\\) : my4\\) uses_allocators\\(memspace\\(\\), traits\\(t2\\) : my3\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : my2\\) uses_allocators\\(memspace\\(3\\), traits\\(t\\) : my\\)" "original" { target c } } } +// { dg-final { scan-tree-dump "#pragma omp target uses_allocators\\(memspace\\(omp_large_cap_mem_space\\), traits\\(\\) : my4\\) uses_allocators\\(memspace\\(\\), traits\\(t2\\) : my3\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : my2\\) uses_allocators\\(memspace\\(omp_high_bw_mem_space\\), traits\\(t\\) : my\\)" "original" { target c++ } } } diff --git a/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-9.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-9.c new file mode 100644 index 000000000000..258dba10ebd8 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-9.c @@ -0,0 +1,44 @@ +#include <stdint.h> +#include <omp.h> + +int main (void) +{ + omp_allocator_handle_t my_alloc; + const omp_alloctrait_t trait[] = { { omp_atk_alignment, 128 } }; + + #pragma omp target uses_allocators (traits (trait): my_alloc) + { + int *x = omp_alloc (sizeof (int), my_alloc); + if ((uintptr_t) x % 128 != 0) + __builtin_abort (); + omp_free (x, my_alloc); + } + + my_alloc = (omp_allocator_handle_t) 0xABCD; + + #pragma omp target uses_allocators (traits (trait): my_alloc) + { + int *x = omp_alloc (sizeof (int), my_alloc); + if ((uintptr_t) x % 128 != 0) + __builtin_abort (); + omp_free (x, my_alloc); + } + + if (my_alloc != (omp_allocator_handle_t) 0xABCD) + __builtin_abort (); + + /* The following creates an allocator with empty traits + default mem space. */ + #pragma omp target uses_allocators (my_alloc) + { + int *x = omp_aligned_alloc (128, sizeof (int), my_alloc); + if ((uintptr_t) x % 128 != 0) + __builtin_abort (); + omp_free (x, my_alloc); + } + + if (my_alloc != (omp_allocator_handle_t) 0xABCD) + __builtin_abort (); + + return 0; +} + diff --git a/libgomp/testsuite/libgomp.c/c.exp b/libgomp/testsuite/libgomp.c/c.exp index aae282478dbc..4b59957d1f3f 100644 --- a/libgomp/testsuite/libgomp.c/c.exp +++ b/libgomp/testsuite/libgomp.c/c.exp @@ -3,6 +3,14 @@ load_gcc_lib gcc-dg.exp lappend ALWAYS_CFLAGS "compiler=$GCC_UNDER_TEST" +proc check_effective_target_c { } { + return 1 +} + +proc check_effective_target_c++ { } { + return 0 +} + # If a testcase doesn't have special options, use these. if ![info exists DEFAULT_CFLAGS] then { set DEFAULT_CFLAGS "-O2" diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators-1.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators-1.f90 new file mode 100644 index 000000000000..cdabbed04e4f --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/uses_allocators-1.f90 @@ -0,0 +1,18 @@ +! { dg-do compile } +use iso_c_binding +use omp_lib +implicit none +contains +subroutine x +integer :: mem +type(omp_alloctrait), parameter:: mem2(1) = [omp_alloctrait(1,1)] +integer(omp_allocator_handle_kind) :: var +!$omp target uses_allocators(memspace(omp_default_mem_space), traits(mem2) : var) defaultmap(none) +block; +type(c_ptr) ::c +c = omp_alloc(omp_default_mem_space, 20_8) +end block +!$omp target uses_allocators(omp_default_mem_alloc, var(mem2)) ! { dg-warning "The specification of arguments to 'uses_allocators' at \\(1\\) where each item is of the form 'allocator\\(traits\\)' is deprecated since OpenMP 5.2; instead use 'uses_allocators\\(traits\\(mem2\\): var\\)' \\\[-Wdeprecated-openmp\\\]" } +block; end block +end +end diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_1.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators-2.f90 similarity index 100% rename from libgomp/testsuite/libgomp.fortran/uses_allocators_1.f90 rename to libgomp/testsuite/libgomp.fortran/uses_allocators-2.f90 diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators-3.f90 similarity index 70% rename from libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90 rename to libgomp/testsuite/libgomp.fortran/uses_allocators-3.f90 index 0ab09975f490..005ead56ee43 100644 --- a/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90 +++ b/libgomp/testsuite/libgomp.fortran/uses_allocators-3.f90 @@ -3,8 +3,6 @@ ! Minimal test for valid code: ! - predefined allocators do not need any special treatment in uses_allocators ! (as 'requires dynamic_allocators' is the default). -! -! - Non-predefined allocators are currently rejected ('sorry)' subroutine test use omp_lib @@ -35,22 +33,22 @@ subroutine non_predef integer(kind=omp_allocator_handle_kind) :: a1, a2, a3 - !$omp target uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2)) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp target uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2)) block; end block - !$omp target parallel uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2)) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp target parallel uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2)) block; end block !$omp target uses_allocators(traits(trait):a1) & - !$omp& uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp& uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3) block; end block !$omp target parallel uses_allocators(traits(trait):a1) & - !$omp& uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp& uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3) block; end block - !$omp target uses_allocators ( traits(trait2) , memspace ( omp_low_lat_mem_space ) : a2 , a3) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp target uses_allocators ( traits(trait2) , memspace ( omp_low_lat_mem_space ) : a2 , a3) block; end block end subroutine @@ -62,7 +60,7 @@ subroutine trait_present integer(kind=omp_allocator_handle_kind) :: a1 ! Invalid in OpenMP 5.0 / 5.1, but valid since 5.2 the same as omp_default_mem_space + emptry traits array - !$omp target uses_allocators ( a1 ) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp target uses_allocators ( a1 ) block; end block end @@ -76,13 +74,13 @@ subroutine odd_names integer(kind=omp_allocator_handle_kind) :: traits integer(kind=omp_allocator_handle_kind) :: memspace - !$omp target uses_allocators ( traits(trait1), memspace(trait1) ) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp target uses_allocators ( traits(trait1), memspace(trait1) ) block; end block - !$omp target uses_allocators ( traits(trait1), memspace(omp_low_lat_mem_space) : traits) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp target uses_allocators ( traits(trait1), memspace(omp_low_lat_mem_space) : traits) block; end block - !$omp target uses_allocators ( memspace(omp_low_lat_mem_space), traits(trait1) : memspace) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp target uses_allocators ( memspace(omp_low_lat_mem_space), traits(trait1) : memspace) block; end block end @@ -94,6 +92,6 @@ subroutine more_checks integer(kind=omp_allocator_handle_kind) :: a1, a2(4) integer(kind=1) :: a3 - !$omp target uses_allocators(memspace (omp_low_lat_mem_space) : a1 ) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp target uses_allocators(memspace (omp_low_lat_mem_space) : a1 ) block; end block end diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators-4.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators-4.f90 new file mode 100644 index 000000000000..0efdc9ef54b0 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/uses_allocators-4.f90 @@ -0,0 +1,59 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } + +program main + use omp_lib + implicit none + integer, allocatable :: arr(:) + integer (omp_allocator_handle_kind) :: bar, foo + + type (omp_alloctrait), parameter :: traits_array(*) = & + [omp_alloctrait(omp_atk_pinned,omp_atv_true),& + omp_alloctrait(omp_atk_partition,omp_atv_nearest)] + + !$omp target allocate(bar : arr) uses_allocators(bar) + block + allocate(arr(100)) + end block + + !$omp target uses_allocators(omp_default_mem_alloc) + block + end block + + !$omp target uses_allocators(bar(traits_array), foo (traits_array)) ! { dg-warning "The specification of arguments to 'uses_allocators' at \\(1\\) where each item is of the form 'allocator\\(traits\\)' is deprecated since OpenMP 5.2; instead use 'uses_allocators\\(traits\\(traits_array\\): foo\\)' \\\[-Wdeprecated-openmp\\\]" } + block + if (foo == 0) stop 1 + end block + + !$omp target uses_allocators(traits(traits_array) : bar) + block + end block + + !$omp target parallel uses_allocators(memspace (omp_low_lat_mem_space) : bar) + block + end block + + !$omp target parallel uses_allocators(memspace (omp_high_bw_mem_space), traits(traits_array) : bar) + block + use iso_c_binding + type(c_ptr) :: ptr + integer(c_size_t) :: sz = 32 + ptr = omp_alloc (sz, bar) + call omp_free (ptr, bar) + end block + +end program main + +! { dg-final { scan-tree-dump "pragma omp target allocate\\(allocator\\(bar\\):arr\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : bar\\)" "original" } } +! { dg-final { scan-tree-dump "pragma omp target" "original" } } +! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(\\), traits\\(traits_array\\) : bar\\) uses_allocators\\(memspace\\(\\), traits\\(traits_array\\) : foo\\)" "original" } } +! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(\\), traits\\(traits_array\\) : bar\\)" "original" } } +! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(omp_low_lat_mem_space\\), traits\\(\\) : bar\\)" "original" } } +! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(omp_high_bw_mem_space\\), traits\\(traits_array\\) : bar\\)" "original" } } + +! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) allocate\\(allocator\\(bar\\):arr\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : bar\\)" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(\\), traits\\(traits_array\\) : bar\\) uses_allocators\\(memspace\\(\\), traits\\(traits_array\\) : foo\\)" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(\\), traits\\(traits_array\\) : bar\\)" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(omp_low_lat_mem_space\\), traits\\(\\) : bar\\)" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(omp_high_bw_mem_space\\), traits\\(traits_array\\) : bar\\)" "gimple" } } diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators-5.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators-5.f90 new file mode 100644 index 000000000000..86a86a08bf0b --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/uses_allocators-5.f90 @@ -0,0 +1,54 @@ +! { dg-do compile } + +program main + use omp_lib + implicit none + integer (omp_allocator_handle_kind) :: bar, foo + + type (omp_alloctrait), parameter :: traits_array(*) = & + [omp_alloctrait(omp_atk_pinned,omp_atv_true),& + omp_alloctrait(omp_atk_partition,omp_atv_nearest)] + + !$omp target uses_allocators(omp_non_existant_alloc) ! { dg-error "Allocator 'omp_non_existant_alloc' at .1. in USES_ALLOCATORS must be a scalar integer of kind 'omp_allocator_handle_kind'" } + block ! { dg-error "Symbol 'omp_non_existant_alloc' at .1. has no IMPLICIT type; did you mean 'omp_const_mem_alloc'\?" "" { target *-*-* } .-1 } + end block + + !$omp target uses_allocators(bar(traits_array), foo (traits_array), ) ! { dg-error "Invalid character in name" } + block + end block + + !$omp target uses_allocators(traits(xyz) : bar) ! { dg-error "Symbol 'xyz' at .1. has no IMPLICIT type" } + block ! { dg-error "Traits array 'xyz' in USES_ALLOCATORS .1. must be a one-dimensional named constant array of type 'omp_alloctrait'" "" { target *-*-* } .-1 } + end block + + !$omp target uses_allocators(memspace(omp_non_existant_mem_space) : foo) ! { dg-error "Symbol 'omp_non_existant_mem_space' at .1. has no IMPLICIT type; did you mean 'omp_const_mem_space'\?" } + ! { dg-error "Memspace 'omp_non_existant_mem_space' at .1. in USES_ALLOCATORS must be a predefined memory space" "" { target *-*-* } .-1 } + + block + end block + + !$omp target uses_allocators(traits(traits_array), traits(traits_array) : bar) ! { dg-error "Duplicate TRAITS modifier at .1. in USES_ALLOCATORS clause" } + block + end block + + !$omp target uses_allocators(memspace(omp_default_mem_space), memspace(omp_default_mem_space) : foo) ! { dg-error "Duplicate MEMSPACE modifier at .1. in USES_ALLOCATORS clause" } + block + end block + + !$omp target uses_allocators(memspace(omp_default_mem_space), traits(traits_array), traits(traits_array) : foo) ! { dg-error "Duplicate TRAITS modifier at .1. in USES_ALLOCATORS clause" } + block + end block + + !$omp target uses_allocators (omp_null_allocator) + block + end block + + !$omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo, bar) + block + end block + + !$omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo(foo_traits)) ! { dg-error "70:Unexpected '\\(' at .1." } + block + end block + +end program main diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators-6.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators-6.f90 new file mode 100644 index 000000000000..00f87109d2c6 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/uses_allocators-6.f90 @@ -0,0 +1,14 @@ +! { dg-do compile } + +program main + use omp_lib + implicit none + integer, allocatable :: arr(:) + integer (omp_allocator_handle_kind) :: bar + + !$omp target allocate(bar : arr) ! { dg-error "allocator 'bar' requires 'uses_allocators.bar.' clause in target region" } + block + allocate(arr(100)) + end block + +end program main diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators-7.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators-7.f90 index 65e2fca082ed..64ac5ca189e4 100644 --- a/libgomp/testsuite/libgomp.fortran/uses_allocators-7.f90 +++ b/libgomp/testsuite/libgomp.fortran/uses_allocators-7.f90 @@ -1,5 +1,4 @@ -! { dg-do compile } -! { dg-additional-options "-fdump-tree-gimple" } +! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } program main use iso_c_binding @@ -28,7 +27,7 @@ program main my_alloc = transfer(int(z'ABCD', omp_allocator_handle_kind), my_alloc) - !$omp target uses_allocators(traits(trait): my_alloc) defaultmap(none) map(tofrom: x, xbuf) + !$omp target uses_allocators(traits(trait): my_alloc) defaultmap(none) map(tofrom: x, xbuf) !$omp parallel allocate(allocator(my_alloc): x, xbuf) if(.false.) firstprivate(x, xbuf) if (mod (TRANSFER (loc(x), iptr), 128) /= 0) & stop 3 @@ -54,16 +53,6 @@ program main stop 8 end - -! FIXME ENABLE: 'dg FIXME final' -> 'dg-final' -! { dg FIXME final { scan-tree-dump-times "#pragma omp target .*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\), traits\\(trait\\)\\)" 1 "gimple" } } -! { dg FIXME final { scan-tree-dump-times "#pragma omp target .*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\), traits\\(\\)\\)" 1 "gimple" } } -! { dg FIXME final { scan-tree-dump "#pragma omp target uses_allocators\\(memspace\\(1\\), traits\\(\\) : my4\\) uses_allocators\\(memspace\\(\\), traits\\(t2\\) : my3\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : my2\\) uses_allocators\\(memspace\\(3\\), traits\\(t\\) : my\\)" 1 "original" } } - - -! FIXME ENABLE code above for "gimple" once it has been implemented: -! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" "" { target *-*-* } 16 } -! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" "" { target *-*-* } 31 } -! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" "" { target *-*-* } 44 } -! { dg-bogus "'my_alloc' not specified in enclosing 'target'" "bogus issue because clause is ignored" { xfail *-*-* } 32 } -! { dg-bogus "'my_alloc' not specified in enclosing 'target'" "bogus issue because clause is ignored" { xfail *-*-* } 45 } +! { dg-final { scan-tree-dump-times "#pragma omp target .*uses_allocators\\(memspace\\(\\), traits\\(trait\\) : my_alloc\\)" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp target .*uses_allocators\\(memspace\\(\\), traits\\(\\) : my_alloc\\)" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp target .*uses_allocators\\(memspace\\(1\\), traits\\(\\) : my4\\) uses_allocators\\(memspace\\(\\), traits\\(t2\\) : my3\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : my2\\) uses_allocators\\(memspace\\(3\\), traits\\(t\\) : my\\)" 1 "original" } } diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators-8.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators-8.f90 new file mode 100644 index 000000000000..63b16d292be7 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/uses_allocators-8.f90 @@ -0,0 +1,51 @@ +program main + use omp_lib + use iso_c_binding + implicit none + integer(c_intptr_t) :: iptr + integer (omp_allocator_handle_kind) :: my_alloc + type(omp_alloctrait), parameter :: trait(*) = & + [omp_alloctrait(omp_atk_alignment, 128)] + + !$omp target uses_allocators(traits(trait) : my_alloc) + block + type(c_ptr) :: ptr + integer(c_size_t) :: sz = 32 + ptr = omp_alloc (sz, my_alloc) + if (mod (TRANSFER (ptr, iptr), 128) /= 0) & + stop 1 + call omp_free (ptr, my_alloc) + end block + + my_alloc = transfer(int(z'ABCD', omp_allocator_handle_kind), my_alloc) + + !$omp target uses_allocators(traits(trait) : my_alloc) + block + type(c_ptr) :: ptr + integer(c_size_t) :: sz = 32 + ptr = omp_alloc (sz, my_alloc) + if (mod (TRANSFER (ptr, iptr), 128) /= 0) & + stop 2 + call omp_free (ptr, my_alloc) + end block + + if (transfer(my_alloc, 0_omp_allocator_handle_kind) /= int(z'ABCD', & + omp_allocator_handle_kind)) & + stop 3 + + ! The following creates an allocator with empty traits + default mem space. + !$omp target uses_allocators(my_alloc) + block + type(c_ptr) :: ptr + integer(c_size_t) :: sz = 32 + ptr = omp_aligned_alloc (128_c_size_t, sz, my_alloc) + if (mod (TRANSFER (ptr, iptr), 128) /= 0) & + stop 4 + call omp_free (ptr, my_alloc) + end block + + if (transfer(my_alloc, 0_omp_allocator_handle_kind) /= int(z'ABCD', & + omp_allocator_handle_kind)) & + stop 5 + +end program main
