This patch is v3 of uses_allocators support.
The issues Tobias raised in the v2 review should all be resolved.
A bit reorg to move the position of MEMSPACE_VALIDATE calls was done.
Largely code logic hasn't really changed.
A new runtime test for alignment had been add for C/C++/Fortran
for a uses_allocators clause + omp_alloc.
Testcases have now all been moved back into libgomp/testsuite,
might consider moving compile tests back to gcc/testsuite after the
"-I<libgomp-build-dir>" patch is approved.
Seeking approval for mainline.
Thanks,
Chung-Lin
2026-06-23 Tobias Burnus <[email protected]>
Andrew Stubbs <[email protected]>
Chung-Lin Tang <[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.
* types.def (BT_FN_VOID_PTRMODE, BT_FN_PTRMODE_PTRMODE,
BT_FN_PTRMODE_PTRMODE_INT_PTR): Add.
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.
* c-c++-common/gomp/uses_allocators-4.c: Likewise.
* gcc.dg/gomp/deprecate-2.c: Likewise.
* gfortran.dg/gomp/allocate-1.f90: Likewise.
* gfortran.dg/gomp/scope-6.f90: Likewise.
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.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 --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 0f96de20074..b2e050141aa 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 9d8e913a84a..d875fea52d4 100644
--- a/gcc/fortran/trans-array.cc
+++ b/gcc/fortran/trans-array.cc
@@ -7213,10 +7213,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)
{
@@ -7287,7 +7284,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 11873d3b6fe..d1d5a7d947b 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -4025,7 +4025,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_;
@@ -5435,14 +5440,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 b1f263a0826..54e245d3ffd 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),
@@ -10868,6 +10899,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:
@@ -10912,6 +10944,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:
@@ -11200,6 +11233,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:
@@ -14826,8 +14860,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:
@@ -14986,6 +15027,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)
{
@@ -15304,6 +15388,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 eef0b418b9b..7ffe65d86e6 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_:
@@ -13215,6 +13254,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)
@@ -13948,6 +13995,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 ce43d70ecf2..00000000000
--- 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 3ab87576c41..00000000000
--- 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 46d286141ef..b7689dd1554 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 8bc6b768778..0463f0e0af9 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 4c4f5e034f7..39a65904c33 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 0674a27501b..96473e52b3b 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -1154,6 +1154,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 def838a949d..3917788d1d1 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 d2e88d97846..5576d92e21d 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 b3d7bcef7d9..f09d19cc755 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 ebbe9816909..577ebb7d1a5 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 87e7855ca3e..c88d212504b 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 a7938beb1ca..6936fce3040 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 60ac7201597..db7c03fb350 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 9510b148ed8..3c8a2a1115e 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 2cc7e6464ea..b0546c87405 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 81426168da9..cc47d205e1e 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 59f2ce07f0e..4ee54e732bd 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -2581,6 +2581,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 47fdd7458ea..cd4f8788e76 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. */
@@ -6256,6 +6307,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 ed096e17b9c..5be949bb611 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 79e51638534..79b5cb0883f 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 df82cbbcba9..2029ee5a9a3 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 c766ea3559d..2babb11018e 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 00000000000..a371d7fe823
--- /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 5f3650157fc..389a51cd1c5 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 00000000000..eeac5f097cb
--- /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 00000000000..23a831e8ba9
--- /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 00000000000..7603dfcd322
--- /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 00000000000..fef13585dbf
--- /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 00000000000..258dba10ebd
--- /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 aae282478db..4b59957d1f3 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 00000000000..cdabbed04e4
--- /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 0ab09975f49..005ead56ee4 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 00000000000..0efdc9ef54b
--- /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 00000000000..86a86a08bf0
--- /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 00000000000..00f87109d2c
--- /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 65e2fca082e..64ac5ca189e 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 00000000000..63b16d292be
--- /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
