This is a new updated patch for OpenMP uses_allocators support.
The last one was submitted by Tobias:
https://gcc.gnu.org/pipermail/gcc-patches/2023-November/637415.html
This new version is a combination of all our patches and fixes during this
period, now including:
1. C/C++ front-end parts re-written, to be more like established style.
2. The target teams issue has been solved by a host-side implementation:
omp_init_allocator() is called on host side, and mapped to target,
with the allocator passed in as a firstprivate variable.
Some adjustments where made to ensure that host/device side must have
same omp_allocator_data format.
3. Various other fixes, e.g. ntraits now using array_type_nelts, don't crash on
VLAs,
omp_null_allocator, etc.
Note that several new tests need my recently submitted testsuite patch to test
correctly:
https://gcc.gnu.org/pipermail/gcc-patches/2025-November/700320.html
Tested without regressions on x86_64-linux + nvptx offloading, amdgcn in
progress but no surprises expected.
Is this okay for mainline?
Thanks,
Chung-Lin
2025-11-12 Tobias Burnus <[email protected]>
Andrew Stubbs <[email protected]>
Chung-Lin Tang <[email protected]>
gcc/ChangeLog:
* builtin-types.def (BT_FN_VOID_PTRMODE): Add.
(BT_FN_PTRMODE_PTRMODE): Add.
(BT_FN_PTRMODE_PTRMODE_INT_PTR): Add.
* gimplify.cc (gimplify_bind_expr): Diagnose missing
uses_allocators clause.
(gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses,
gimplify_omp_workshare): Handle uses_allocators.
* omp-builtins.def (BUILT_IN_OMP_INIT_ALLOCATOR): Add.
(BUILT_IN_OMP_DESTROY_ALLOCATOR): Add.
(BUILT_IN_GOMP_OMP_ALLOCATOR_MAP): Add.
(BUILT_IN_GOMP_OMP_ALLOCATOR_UNMAP): Add.
* omp-low.cc (scan_sharing_clauses): Check if ALLOCATE clause in target
region has containing USES_ALLOCATORS clause.
Add OMP_CLAUSE_USES_ALLOCATORS cases.
* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_USES_ALLOCATORS.
* tree.cc (omp_clause_num_ops, omp_clause_code_name): Likewise.
* tree-pretty-print.cc (dump_omp_clause): Handle it.
* tree.h (OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR,
OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE,
OMP_CLAUSE_USES_ALLOCATORS_TRAITS): New.
gcc/c-family/ChangeLog:
* c-omp.cc (c_omp_split_clauses): Hande uses_allocators.
* c-pragma.h (enum pragma_omp_clause): Add
PRAGMA_OMP_CLAUSE_USES_ALLOCATORS.
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_clause_uses_allocators): New function.
(c_parser_omp_clause_name, c_parser_omp_all_clauses,
OMP_TARGET_CLAUSE_MASK): Handle uses_allocators.
* c-typeck.cc (c_finish_omp_clauses): Likewise.
gcc/cp/ChangeLog:
* parser.cc (cp_parser_omp_clause_uses_allocators): New function.
(cp_parser_omp_clause_name, cp_parser_omp_all_clauses,
OMP_TARGET_CLAUSE_MASK): Handle uses_allocators.
* semantics.cc (finish_omp_clauses): Likewise.
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.
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_omp_allocator_data_size): New variable for exporting size of
struct omp_allocator_data within libgomp.
(gomp_memspace_validate): New function, enabled only for host-side.
* 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): Remove.
* configure.ac: Require offload hosts and accelerator targets to support
__sync_* builtins.
* configure: Regenerate.
* libgomp-plugin.h (nvptx_memspace_validate): New declaration.
* libgomp.h (gomp_omp_allocator_data_size): New declaration.
(struct gomp_device_descr): New memspace_validate_func hook.
* libgomp.map (GOMP_6.0.1): Add GOMP_omp_allocator_map and
GOMP_omp_allocator_unmap.
* oacc-host.c (host_dispatch): Init memspace_validate_func field.
* plugin/configfrag.ac: Define OFFLOAD_PLUGINS only if none.
* plugin/plugin-gcn.c (GOMP_OFFLOAD_memspace_validate): New function.
* plugin/plugin-nvptx.c (GOMP_OFFLOAD_memspace_validate): New function.
* target.c (GOMP_omp_allocator_map): New function.
(GOMP_omp_allocator_unmap): New function.
(gomp_load_plugin_for_device): Add DLSYM_OPT of memspace_validate_func
hook.
* 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.
* testsuite/libgomp.fortran/uses_allocators-7.f90: New test.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/uses_allocators-1.c: New test.
* c-c++-common/gomp/uses_allocators-2.c: New test.
* c-c++-common/gomp/uses_allocators-3.c: New test.
* c-c++-common/gomp/uses_allocators-4.c: New test.
* c-c++-common/gomp/uses_allocators-5.c: New test.
* c-c++-common/gomp/uses_allocators-6.c: New test.
* gfortran.dg/gomp/allocate-1.f90: Add uses_allocators.
* gfortran.dg/gomp/scope-6.f90: Update dg-scan-tree-dump.
* gfortran.dg/gomp/uses_allocators-1.f90: New test.
* gfortran.dg/gomp/uses_allocators-2.f90: New test.
* gfortran.dg/gomp/uses_allocators-3.f90: New test.
* gfortran.dg/gomp/uses_allocators-4.f90: New test.
* gfortran.dg/gomp/uses_allocators-5.f90: New test.
* gfortran.dg/gomp/uses_allocators-6.f90: New test.
* gfortran.dg/gomp/uses_allocators-7.f90: New test.
diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index 9583d30dfc0..ad166a64091 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -393,6 +393,7 @@ DEF_FUNCTION_TYPE_1 (BT_FN_DFLOAT64_DFLOAT64, BT_DFLOAT64,
BT_DFLOAT64)
DEF_FUNCTION_TYPE_1 (BT_FN_DFLOAT128_DFLOAT128, BT_DFLOAT128, BT_DFLOAT128)
DEF_FUNCTION_TYPE_1 (BT_FN_DFLOAT64X_DFLOAT64X, BT_DFLOAT64X, BT_DFLOAT64X)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_VPTR, BT_VOID, BT_VOLATILE_PTR)
+DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRMODE, BT_VOID, BT_PTRMODE)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRPTR, BT_VOID, BT_PTR_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_CONST_PTR, BT_VOID, BT_CONST_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_UINT_UINT, BT_UINT, BT_UINT)
@@ -420,6 +421,7 @@ DEF_FUNCTION_TYPE_1 (BT_FN_UINT64_FLOAT, BT_UINT64,
BT_FLOAT)
DEF_FUNCTION_TYPE_1 (BT_FN_BOOL_INT, BT_BOOL, BT_INT)
DEF_FUNCTION_TYPE_1 (BT_FN_BOOL_PTR, BT_BOOL, BT_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_PTR_CONST_PTR, BT_PTR, BT_CONST_PTR)
+DEF_FUNCTION_TYPE_1 (BT_FN_PTRMODE_PTRMODE, BT_PTRMODE, BT_PTRMODE)
DEF_FUNCTION_TYPE_1 (BT_FN_CONST_PTR_CONST_PTR, BT_CONST_PTR, BT_CONST_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_UINT16_UINT32, BT_UINT16, BT_UINT32)
DEF_FUNCTION_TYPE_1 (BT_FN_UINT32_UINT16, BT_UINT32, BT_UINT16)
@@ -862,6 +864,8 @@ DEF_FUNCTION_TYPE_3 (BT_FN_UINT64_UINT64_UINT32_CONST_SIZE,
BT_UINT64,
BT_UINT64, BT_UINT32, BT_CONST_SIZE)
DEF_FUNCTION_TYPE_3 (BT_FN_UINT64_UINT64_UINT64_CONST_SIZE, BT_UINT64,
BT_UINT64, BT_UINT64, BT_CONST_SIZE)
+DEF_FUNCTION_TYPE_3 (BT_FN_PTRMODE_PTRMODE_INT_PTR, BT_PTRMODE, BT_PTRMODE,
+ BT_INT, BT_PTR)
DEF_FUNCTION_TYPE_4 (BT_FN_SIZE_CONST_PTR_SIZE_SIZE_FILEPTR,
BT_SIZE, BT_CONST_PTR, BT_SIZE, BT_SIZE, BT_FILEPTR)
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 3c2ee9ff1ae..fe8c6456832 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -2181,6 +2181,7 @@ c_omp_split_clauses (location_t loc, enum tree_code code,
case OMP_CLAUSE_HAS_DEVICE_ADDR:
case OMP_CLAUSE_DEFAULTMAP:
case OMP_CLAUSE_DEPEND:
+ case OMP_CLAUSE_USES_ALLOCATORS:
s = C_OMP_CLAUSE_SPLIT_TARGET;
break;
case OMP_CLAUSE_DOACROSS:
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 13df9ea490e..21cbc34e8a8 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -174,6 +174,7 @@ enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_USE,
PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR,
PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR,
+ PRAGMA_OMP_CLAUSE_USES_ALLOCATORS,
/* Clauses for OpenACC. */
PRAGMA_OACC_CLAUSE_ASYNC,
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 360df52967b..0e58541db56 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -16418,6 +16418,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR;
else if (!strcmp ("use_device_ptr", p))
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
+ else if (!strcmp ("uses_allocators", p))
+ result = PRAGMA_OMP_CLAUSE_USES_ALLOCATORS;
break;
case 'v':
if (!strcmp ("vector", p))
@@ -19386,6 +19388,220 @@ c_parser_omp_clause_allocate (c_parser *parser, tree
list)
return nl;
}
+/* OpenMP 5.0:
+ uses_allocators ( allocator-list )
+
+ allocator-list:
+ allocator
+ allocator , allocator-list
+ allocator ( traits-array )
+ allocator ( traits-array ) , allocator-list
+
+ OpenMP 5.2:
+
+ uses_allocators ( modifier : allocator-list )
+ uses_allocators ( modifier , modifier : allocator-list )
+
+ modifier:
+ traits ( traits-array )
+ memspace ( mem-space-handle ) */
+
+static tree
+c_parser_omp_clause_uses_allocators (c_parser *parser, tree list)
+{
+ location_t clause_loc = c_parser_peek_token (parser)->location;
+ tree nl = list;
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
+
+ bool has_modifiers = false;
+ bool seen_allocators = false;
+ tree memspace_expr = NULL_TREE;
+ tree traits_var = NULL_TREE;
+
+ if (c_parser_next_token_is (parser, CPP_NAME)
+ && c_parser_peek_2nd_token (parser)->type == CPP_OPEN_PAREN)
+ {
+ unsigned int n = 3;
+ const char *p
+ = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+ if ((strcmp (p, "traits") == 0 || strcmp (p, "memspace") == 0)
+ && c_parser_check_balanced_raw_token_sequence (parser, &n)
+ && (c_parser_peek_nth_token_raw (parser, n)->type
+ == CPP_CLOSE_PAREN))
+ {
+ if (c_parser_peek_nth_token_raw (parser, n + 1)->type
+ == CPP_COLON)
+ has_modifiers = true;
+ else if (c_parser_peek_nth_token_raw (parser, n + 1)->type
+ == CPP_COMMA
+ && (c_parser_peek_nth_token_raw (parser, n + 2)->type
+ == CPP_NAME)
+ && (c_parser_peek_nth_token_raw (parser, n + 3)->type
+ == CPP_OPEN_PAREN))
+ {
+ c_token *tok = c_parser_peek_nth_token_raw (parser, n + 2);
+ const char *q = IDENTIFIER_POINTER (tok->value);
+ n += 4;
+ if ((strcmp (q, "traits") == 0
+ || strcmp (q, "memspace") == 0)
+ && c_parser_check_balanced_raw_token_sequence (parser, &n)
+ && (c_parser_peek_nth_token_raw (parser, n)->type
+ == CPP_CLOSE_PAREN))
+ {
+ if (c_parser_peek_nth_token_raw (parser, n + 1)->type
+ == CPP_COLON)
+ has_modifiers = true;
+ if ((c_parser_peek_nth_token_raw (parser, n + 1)->type
+ == CPP_COMMA)
+ && (c_parser_peek_nth_token_raw (parser, n + 2)->type
+ == CPP_NAME))
+ {
+ c_token *tok
+ = c_parser_peek_nth_token_raw (parser, n + 2);
+ const char *m = IDENTIFIER_POINTER (tok->value);
+ if (strcmp (p, m) == 0 || strcmp (q, m) == 0)
+ {
+ error_at (tok->location, "duplicate %qs modifier", m);
+ goto end;
+ }
+ }
+ }
+ }
+ }
+ if (has_modifiers)
+ {
+ c_parser_consume_token (parser);
+ matching_parens parens2;
+ parens2.require_open (parser);
+ c_expr expr = c_parser_expr_no_commas (parser, NULL);
+ if (expr.value == error_mark_node)
+ ;
+ else if (strcmp (p, "traits") == 0)
+ {
+ traits_var = expr.value;
+ traits_var = c_fully_fold (traits_var, false, NULL);
+ }
+ else
+ {
+ memspace_expr = expr.value;
+ memspace_expr = c_fully_fold (memspace_expr, false, NULL);
+ }
+ parens2.skip_until_found_close (parser);
+ if (c_parser_next_token_is (parser, CPP_COMMA))
+ {
+ c_parser_consume_token (parser);
+ c_token *tok = c_parser_peek_token (parser);
+ const char *q = "";
+ if (c_parser_next_token_is (parser, CPP_NAME))
+ q = IDENTIFIER_POINTER (tok->value);
+ if (strcmp (q, "traits") != 0 && strcmp (q, "memspace") != 0)
+ {
+ c_parser_error (parser, "expected %<traits%> or "
+ "%<memspace%>");
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ else if (strcmp (p, q) == 0)
+ {
+ error_at (tok->location, "duplicate %qs modifier", p);
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ c_parser_consume_token (parser);
+ if (!parens2.require_open (parser))
+ {
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ expr = c_parser_expr_no_commas (parser, NULL);
+ if (strcmp (q, "traits") == 0)
+ {
+ traits_var = expr.value;
+ traits_var = c_fully_fold (traits_var, false, NULL);
+ }
+ else
+ {
+ memspace_expr = expr.value;
+ memspace_expr = c_fully_fold (memspace_expr, false, NULL);
+ }
+ parens2.skip_until_found_close (parser);
+ }
+ if (!c_parser_require (parser, CPP_COLON, "expected %<:%>"))
+ goto end;
+ }
+ }
+
+ while (c_parser_next_token_is (parser, CPP_NAME))
+ {
+ c_token *tok = c_parser_peek_token (parser);
+ const char *tok_s = IDENTIFIER_POINTER (tok->value);
+ tree t = lookup_name (tok->value);
+ if (t == NULL_TREE)
+ {
+ undeclared_variable (tok->location, tok->value);
+ t = error_mark_node;
+ }
+ c_parser_consume_token (parser);
+
+ /* Legacy traits syntax. */
+ tree legacy_traits = NULL_TREE;
+ if (c_parser_next_token_is (parser, CPP_OPEN_PAREN)
+ && c_parser_peek_2nd_token (parser)->type == CPP_NAME
+ && c_parser_peek_nth_token_raw (parser, 3)->type == CPP_CLOSE_PAREN)
+ {
+ matching_parens parens2;
+ parens2.require_open (parser);
+ const char *tok_a
+ = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+ location_t expr_loc = c_parser_peek_token (parser)->location;
+ c_expr expr = c_parser_expr_no_commas (parser, NULL);
+ parens2.skip_until_found_close (parser);
+
+ if (has_modifiers)
+ {
+ error_at (expr_loc,
+ "legacy %<%s(%s)%> traits syntax not allowed in "
+ "%<uses_allocators%> clause when using modifiers",
+ tok_s, tok_a);
+ goto end;
+ }
+
+ legacy_traits = c_fully_fold (expr.value, false, NULL);
+ if (legacy_traits == error_mark_node)
+ goto end;
+ }
+
+ if (seen_allocators && has_modifiers)
+ {
+ error_at (c_parser_peek_token (parser)->location,
+ "%<uses_allocators%> clause only accepts a single "
+ "allocator when using modifiers");
+ goto end;
+ }
+ seen_allocators = true;
+
+ tree c = build_omp_clause (clause_loc,
+ OMP_CLAUSE_USES_ALLOCATORS);
+ OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c) = t;
+ OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) = memspace_expr;
+ OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c) = (legacy_traits
+ ? legacy_traits : traits_var);
+ OMP_CLAUSE_CHAIN (c) = nl;
+ nl = c;
+
+ if (c_parser_next_token_is (parser, CPP_COMMA))
+ c_parser_consume_token (parser);
+ else
+ break;
+ }
+
+ end:
+ parens.skip_until_found_close (parser);
+ return nl;
+}
+
/* OpenMP 4.0:
linear ( variable-list )
linear ( variable-list : expression )
@@ -21797,6 +22013,10 @@ c_parser_omp_all_clauses (c_parser *parser,
omp_clause_mask mask,
clauses = c_parser_omp_clause_linear (parser, clauses);
c_name = "linear";
break;
+ case PRAGMA_OMP_CLAUSE_USES_ALLOCATORS:
+ clauses = c_parser_omp_clause_uses_allocators (parser, clauses);
+ c_name = "uses_allocators";
+ break;
case PRAGMA_OMP_CLAUSE_AFFINITY:
clauses = c_parser_omp_clause_affinity (parser, clauses);
c_name = "affinity";
@@ -26638,8 +26858,9 @@ c_parser_omp_target_exit_data (location_t loc, c_parser
*parser,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)\
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR))
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USES_ALLOCATORS))
static bool
c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index e1d2d1173dc..aa90288f81b 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -17205,6 +17205,128 @@ c_finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
break;
}
gcc_unreachable ();
+
+ case OMP_CLAUSE_USES_ALLOCATORS:
+ t = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
+ if (t == error_mark_node)
+ {
+ remove = true;
+ break;
+ }
+ if ((VAR_P (t) || TREE_CODE (t) == PARM_DECL)
+ && (bitmap_bit_p (&generic_head, DECL_UID (t))
+ || bitmap_bit_p (&map_head, DECL_UID (t))
+ || bitmap_bit_p (&firstprivate_head, DECL_UID (t))
+ || bitmap_bit_p (&lastprivate_head, DECL_UID (t))))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE appears more than once in data clauses", t);
+ remove = true;
+ break;
+ }
+ else
+ bitmap_set_bit (&generic_head, DECL_UID (t));
+ if (TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE
+ || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))),
+ "omp_allocator_handle_t") != 0)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "allocator must be of %<omp_allocator_handle_t%> type");
+ remove = true;
+ break;
+ }
+ if (TREE_CODE (t) == CONST_DECL)
+ {
+ /* Currently for pre-defined allocators in libgomp, we do not
+ require additional init/fini inside target regions, so discard
+ such clauses. */
+ remove = true;
+
+ if (OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c)
+ || OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "modifiers cannot be used with pre-defined "
+ "allocators");
+ break;
+ }
+ }
+ t = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c);
+ if (t == error_mark_node)
+ {
+ remove = true;
+ break;
+ }
+ if (t != NULL_TREE
+ && ((TREE_CODE (t) != CONST_DECL && TREE_CODE (t) != INTEGER_CST)
+ || TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE
+ || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE
(t))),
+ "omp_memspace_handle_t") != 0))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "memspace modifier must be "
+ "constant enum of %<omp_memspace_handle_t%> type");
+ remove = true;
+ break;
+ }
+ t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);
+ if (t == error_mark_node)
+ {
+ remove = true;
+ break;
+ }
+ if (t != NULL_TREE
+ && t != error_mark_node
+ && (DECL_EXTERNAL (t)
+ || TREE_CODE (t) == PARM_DECL))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "traits array %qE must be "
+ "defined in same scope as the construct on which the "
+ "clause appears", t);
+ remove = true;
+ }
+ if (t != NULL_TREE)
+ {
+ bool type_err = false;
+
+ if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE
+ || DECL_SIZE (t) == NULL_TREE
+ || !COMPLETE_TYPE_P (TREE_TYPE (t)))
+ type_err = true;
+ else
+ {
+ tree elem_t = TREE_TYPE (TREE_TYPE (t));
+ if (TREE_CODE (elem_t) != RECORD_TYPE
+ || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (elem_t)),
+ "omp_alloctrait_t") != 0
+ || !TYPE_READONLY (elem_t))
+ type_err = true;
+ }
+ if (type_err)
+ {
+ if (t != error_mark_node)
+ error_at (OMP_CLAUSE_LOCATION (c), "traits array %qE must "
+ "be of %<const omp_alloctrait_t []%> type", t);
+ else
+ error_at (OMP_CLAUSE_LOCATION (c), "traits array must "
+ "be of %<const omp_alloctrait_t []%> type");
+ remove = true;
+ }
+ else
+ {
+ tree cst_val = decl_constant_value_1 (t, true);
+ if (cst_val == t)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "traits array must be "
+ "initialized with constants");
+
+ remove = true;
+ }
+ }
+ }
+ if (remove)
+ break;
+ pc = &OMP_CLAUSE_CHAIN (c);
+ continue;
case OMP_CLAUSE_DEPEND:
depend_clause = c;
/* FALLTHRU */
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 06cba31ada6..4e3146dd7be 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -39765,6 +39765,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR;
else if (!strcmp ("use_device_ptr", p))
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
+ else if (!strcmp ("uses_allocators", p))
+ result = PRAGMA_OMP_CLAUSE_USES_ALLOCATORS;
break;
case 'v':
if (!strcmp ("vector", p))
@@ -42317,6 +42319,185 @@ cp_parser_omp_clause_allocate (cp_parser *parser,
tree list)
return nlist;
}
+/* OpenMP 5.0:
+ uses_allocators ( allocator-list )
+
+ allocator-list:
+ allocator
+ allocator , allocator-list
+ allocator ( traits-array )
+ allocator ( traits-array ) , allocator-list
+
+ OpenMP 5.2:
+
+ uses_allocators ( modifier : allocator-list )
+ uses_allocators ( modifier , modifier : allocator-list )
+
+ modifier:
+ traits ( traits-array )
+ memspace ( mem-space-handle ) */
+
+static tree
+cp_parser_omp_clause_uses_allocators (cp_parser *parser, tree list)
+{
+ location_t clause_loc
+ = cp_lexer_peek_token (parser->lexer)->location;
+ tree nl = list;
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
+
+ bool has_modifiers = false;
+ bool seen_allocators = false;
+ tree memspace_expr = NULL_TREE;
+ tree traits_var = NULL_TREE;
+
+ cp_parser_parse_tentatively (parser);
+ bool saved_colon_corrects_to_scope_p = parser->colon_corrects_to_scope_p;
+ parser->colon_corrects_to_scope_p = false;
+
+ cp_token *dup_mod_tok = NULL;
+ for (int mod = 0; mod <= 2; mod++)
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)
+ && cp_lexer_nth_token_is (parser->lexer, 2, CPP_OPEN_PAREN))
+ {
+ cp_token *mod_tok = cp_lexer_peek_token (parser->lexer);
+ tree id = mod_tok->u.value;
+ const char *p = IDENTIFIER_POINTER (id);
+ if (strcmp (p, "traits") != 0 && strcmp (p, "memspace") != 0)
+ break;
+ cp_lexer_consume_token (parser->lexer);
+ matching_parens parens2;
+ if (!parens2.require_open (parser))
+ break;
+ tree t = cp_parser_assignment_expression (parser);
+ if (strcmp (p, "traits") == 0)
+ {
+ if (traits_var != NULL_TREE)
+ dup_mod_tok = mod_tok;
+ else
+ traits_var = t;
+ }
+ else
+ {
+ if (memspace_expr != NULL_TREE)
+ dup_mod_tok = mod_tok;
+ else
+ memspace_expr = t;
+ }
+ if (!parens2.require_close (parser))
+ break;
+ if (cp_lexer_next_token_is (parser->lexer, CPP_COLON))
+ {
+ has_modifiers = true;
+ cp_lexer_consume_token (parser->lexer);
+ break;
+ }
+ if (/*mod != 0 || */ cp_lexer_next_token_is_not (parser->lexer,
CPP_COMMA))
+ break;
+ cp_lexer_consume_token (parser->lexer);
+ }
+ else
+ break;
+
+ if (!has_modifiers)
+ {
+ cp_parser_abort_tentative_parse (parser);
+ traits_var = NULL_TREE;
+ memspace_expr = NULL_TREE;
+ }
+ else
+ {
+ if (dup_mod_tok)
+ {
+ error_at (dup_mod_tok->location, "duplicate %qs modifier",
+ IDENTIFIER_POINTER (dup_mod_tok->u.value));
+ cp_parser_parse_definitely (parser);
+ goto end;
+ }
+ cp_parser_parse_definitely (parser);
+ }
+ parser->colon_corrects_to_scope_p = saved_colon_corrects_to_scope_p;
+
+ while (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ {
+ cp_token *tok = cp_lexer_peek_token (parser->lexer);
+ tree t;
+ t = cp_parser_lookup_name_simple (parser,
+ tok->u.value,
+ tok->location);
+ if (t == error_mark_node)
+ cp_parser_name_lookup_error (parser, tok->u.value, t, NLE_NULL,
+ tok->location);
+ cp_lexer_consume_token (parser->lexer);
+
+ /* Legacy traits syntax. */
+ tree legacy_traits = NULL_TREE;
+ if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_PAREN)
+ && cp_lexer_nth_token_is (parser->lexer, 2, CPP_NAME)
+ && cp_lexer_nth_token_is (parser->lexer, 3, CPP_CLOSE_PAREN))
+ {
+ matching_parens parens2;
+ parens2.require_open (parser);
+ cp_token *arg_tok = cp_lexer_peek_token (parser->lexer);
+
+ tree arg = cp_parser_lookup_name_simple (parser, arg_tok->u.value,
+ arg_tok->location);
+ if (arg == error_mark_node)
+ cp_parser_name_lookup_error (parser, arg_tok->u.value, arg,
+ NLE_NULL, arg_tok->location);
+ cp_lexer_consume_token (parser->lexer);
+ parens2.require_close (parser);
+
+ if (has_modifiers)
+ {
+ error_at (arg_tok->location,
+ "legacy %<%E(%E)%> traits syntax not allowed in "
+ "%<uses_allocators%> clause when using modifiers",
+ tok->u.value, arg_tok->u.value);
+ goto end;
+ }
+
+ legacy_traits = arg;
+ if (legacy_traits == error_mark_node)
+ goto end;
+ }
+
+ if (seen_allocators && has_modifiers)
+ {
+ error_at (cp_lexer_peek_token (parser->lexer)->location,
+ "%<uses_allocators%> clause only accepts a single "
+ "allocator when using modifiers");
+ goto end;
+ }
+ seen_allocators = true;
+
+ tree c = build_omp_clause (clause_loc,
+ OMP_CLAUSE_USES_ALLOCATORS);
+ OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c) = t;
+ OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) = memspace_expr;
+ OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c) = (legacy_traits
+ ? legacy_traits : traits_var);
+ OMP_CLAUSE_CHAIN (c) = nl;
+ nl = c;
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
+ cp_lexer_consume_token (parser->lexer);
+ else
+ break;
+ }
+
+ if (!parens.require_close (parser))
+ goto end;
+ return nl;
+ end:
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return nl;
+}
+
/* OpenMP 2.5:
lastprivate ( variable-list )
@@ -44870,6 +45051,10 @@ cp_parser_omp_all_clauses (cp_parser *parser,
omp_clause_mask mask,
clauses = cp_parser_omp_clause_allocate (parser, clauses);
c_name = "allocate";
break;
+ case PRAGMA_OMP_CLAUSE_USES_ALLOCATORS:
+ clauses = cp_parser_omp_clause_uses_allocators (parser, clauses);
+ c_name = "uses_allocators";
+ break;
case PRAGMA_OMP_CLAUSE_LINEAR:
{
bool declare_simd = false;
@@ -49788,7 +49973,8 @@ cp_parser_omp_target_update (cp_parser *parser,
cp_token *pragma_tok,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)\
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR))
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR)\
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USES_ALLOCATORS))
static bool
cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 3e19a56f51e..053b3ae2173 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -8870,6 +8870,125 @@ finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
break;
}
gcc_unreachable ();
+ case OMP_CLAUSE_USES_ALLOCATORS:
+ t = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
+ if (t == error_mark_node)
+ {
+ remove = true;
+ break;
+ }
+ if (TREE_CODE (t) == FIELD_DECL)
+ {
+ sorry_at (OMP_CLAUSE_LOCATION (c), "class members not yet "
+ "supported in %<uses_allocators%> clause");
+ remove = true;
+ break;
+ }
+ t = convert_from_reference (t);
+ if (TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE
+ || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))),
+ "omp_allocator_handle_t") != 0)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "allocator must be of %<omp_allocator_handle_t%> type");
+ remove = true;
+ break;
+ }
+ if (TREE_CODE (t) == CONST_DECL)
+ {
+ /* Currently for pre-defined allocators in libgomp, we do not
+ require additional init/fini inside target regions, so discard
+ such clauses. */
+ remove = true;
+
+ if (OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c)
+ || OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "modifiers cannot be used with pre-defined "
+ "allocators");
+ break;
+ }
+ }
+ t = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c);
+ if (t == error_mark_node)
+ {
+ remove = true;
+ break;
+ }
+ if (t != NULL_TREE
+ && ((TREE_CODE (t) != CONST_DECL && TREE_CODE (t) != INTEGER_CST)
+ || TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE
+ || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE
(t))),
+ "omp_memspace_handle_t") != 0))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "memspace modifier must be "
+ "constant enum of %<omp_memspace_handle_t%> type");
+ remove = true;
+ break;
+ }
+ t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);
+ if (t == error_mark_node)
+ {
+ remove = true;
+ break;
+ }
+ if (t != NULL_TREE
+ && t != error_mark_node
+ && (DECL_EXTERNAL (t)
+ || TREE_CODE (t) == PARM_DECL))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "traits array %qE must be "
+ "defined in same scope as the construct on which the "
+ "clause appears", t);
+ remove = true;
+ }
+ if (t != NULL_TREE)
+ {
+ bool type_err = false;
+
+ if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE
+ || DECL_SIZE (t) == NULL_TREE
+ || !COMPLETE_TYPE_P (TREE_TYPE (t)))
+ type_err = true;
+ else
+ {
+ tree elem_t = TREE_TYPE (TREE_TYPE (t));
+ if (TREE_CODE (elem_t) != RECORD_TYPE
+ || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (elem_t)),
+ "omp_alloctrait_t") != 0
+ || !TYPE_READONLY (elem_t))
+ type_err = true;
+ }
+ if (type_err)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "traits array %qE must "
+ "be of %<const omp_alloctrait_t []%> type", t);
+ remove = true;
+ }
+ else if (TREE_CODE (array_type_nelts_top (TREE_TYPE (t)))
+ != INTEGER_CST)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "variable length traits "
+ "arrays are not supported");
+ remove = true;
+ }
+ else
+ {
+ tree cst_val = decl_constant_value (t);
+ if (cst_val == t)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "traits array must be "
+ "initialized with constants");
+
+ remove = true;
+ }
+ }
+ }
+ if (remove)
+ break;
+ pc = &OMP_CLAUSE_CHAIN (c);
+ continue;
case OMP_CLAUSE_DEPEND:
depend_clause = c;
/* FALLTHRU */
diff --git a/gcc/fortran/trans-array.cc b/gcc/fortran/trans-array.cc
index cd137212260..ee4c4a3a736 100644
--- a/gcc/fortran/trans-array.cc
+++ b/gcc/fortran/trans-array.cc
@@ -7078,10 +7078,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)
{
@@ -7152,7 +7149,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 69a70d7138c..e88a3d0255d 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -3728,7 +3728,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_;
@@ -5091,13 +5096,36 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
}
break;
case OMP_LIST_USES_ALLOCATORS:
- /* Ignore 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/fortran/types.def b/gcc/fortran/types.def
index dd9b8df59be..51b5ef6bcc8 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -80,11 +80,13 @@ DEF_FUNCTION_TYPE_0 (BT_FN_UINT, BT_UINT)
DEF_FUNCTION_TYPE_0 (BT_FN_VOID, BT_VOID)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTR, BT_VOID, BT_PTR)
+DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRMODE, BT_VOID, BT_PTRMODE)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRPTR, BT_VOID, BT_PTR_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_VPTR, BT_VOID, BT_VOLATILE_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_INT_INT, BT_INT, BT_INT)
DEF_FUNCTION_TYPE_1 (BT_FN_UINT_UINT, BT_UINT, BT_UINT)
DEF_FUNCTION_TYPE_1 (BT_FN_PTR_PTR, BT_PTR, BT_PTR)
+DEF_FUNCTION_TYPE_1 (BT_FN_PTRMODE_PTRMODE, BT_PTRMODE, BT_PTRMODE)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_INT, BT_VOID, BT_INT)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_BOOL, BT_VOID, BT_BOOL)
DEF_FUNCTION_TYPE_1 (BT_FN_BOOL_INT, BT_BOOL, BT_INT)
@@ -156,6 +158,8 @@ DEF_FUNCTION_TYPE_3 (BT_FN_VOID_SIZE_SIZE_PTR, BT_VOID,
BT_SIZE, BT_SIZE,
DEF_FUNCTION_TYPE_3 (BT_FN_UINT_UINT_PTR_PTR, BT_UINT, BT_UINT, BT_PTR, BT_PTR)
DEF_FUNCTION_TYPE_3 (BT_FN_PTR_SIZE_SIZE_PTRMODE,
BT_PTR, BT_SIZE, BT_SIZE, BT_PTRMODE)
+DEF_FUNCTION_TYPE_3 (BT_FN_PTRMODE_PTRMODE_INT_PTR, BT_PTRMODE, BT_PTRMODE,
+ BT_INT, BT_PTR)
DEF_FUNCTION_TYPE_4 (BT_FN_PTR_PTR_SIZE_PTRMODE_PTRMODE,
BT_PTR, BT_PTR, BT_SIZE, BT_PTRMODE, BT_PTRMODE)
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index d8725e4c5e2..297db983411 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -1449,18 +1449,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),
@@ -14832,6 +14860,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
nowait = 1;
break;
+ case OMP_CLAUSE_USES_ALLOCATORS:
+ 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_FIRSTPRIVATE);
+
+ decl = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);
+ if (decl && !DECL_INITIAL (decl))
+ omp_add_variable (ctx, decl, GOVD_SEEN | GOVD_FIRSTPRIVATE);
+ }
+ else
+ remove = true;
+ break;
+
case OMP_CLAUSE_ORDERED:
case OMP_CLAUSE_UNTIED:
case OMP_CLAUSE_COLLAPSE:
@@ -14982,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)
{
@@ -16335,6 +16421,7 @@ end_adjust_omp_map_clause:
case OMP_CLAUSE_FINALIZE:
case OMP_CLAUSE_INCLUSIVE:
case OMP_CLAUSE_EXCLUSIVE:
+ case OMP_CLAUSE_USES_ALLOCATORS:
break;
case OMP_CLAUSE_NOHOST:
@@ -18759,6 +18846,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
tree expr = *expr_p;
gimple *stmt;
gimple_seq body = NULL;
+ gimple_seq pre_stmt_seq = NULL, post_stmt_seq = NULL;
enum omp_region_type ort;
switch (TREE_CODE (expr))
@@ -18847,6 +18935,107 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq
*pre_p)
body = NULL;
gimple_seq_add_stmt (&body, g);
}
+ else if ((ort & ORT_TARGET) != 0 && (ort & ORT_ACC) == 0)
+ {
+ gimple_seq init_seq = NULL;
+ gimple_seq fini_seq = NULL;
+
+ tree omp_init_allocator_fn = NULL_TREE;
+ tree omp_destroy_allocator_fn = NULL_TREE;
+
+ for (tree *cp = &OMP_CLAUSES (expr); *cp != NULL;
+ cp = &OMP_CLAUSE_CHAIN (*cp))
+ if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_USES_ALLOCATORS)
+ {
+ tree c = *cp;
+ tree allocator = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
+ tree memspace = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c);
+ tree traits = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);
+
+ if (omp_init_allocator_fn == NULL_TREE)
+ {
+ omp_init_allocator_fn
+ = builtin_decl_explicit (BUILT_IN_OMP_INIT_ALLOCATOR);
+ omp_destroy_allocator_fn
+ = builtin_decl_explicit (BUILT_IN_OMP_DESTROY_ALLOCATOR);
+ }
+ 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, &pre_stmt_seq, 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 omp_allocator_map_fn
+ = builtin_decl_explicit (BUILT_IN_GOMP_OMP_ALLOCATOR_MAP);
+ tree omp_allocator_unmap_fn
+ = builtin_decl_explicit (BUILT_IN_GOMP_OMP_ALLOCATOR_UNMAP);
+
+ tree host_allocator = create_tmp_var (TREE_TYPE (allocator),
+ "host_allocator");
+ tree orig_allocator = create_tmp_var (TREE_TYPE (allocator),
+ "orig_allocator");
+
+ tree call = build_call_expr_loc (OMP_CLAUSE_LOCATION (c),
+ omp_init_allocator_fn, 3,
+ memspace, ntraits,
+ traits_var);
+ call = fold_convert (TREE_TYPE (allocator), call);
+ gimplify_assign (host_allocator, call, &pre_stmt_seq);
+
+ tree map_call = build_call_expr_loc (OMP_CLAUSE_LOCATION (c),
+ omp_allocator_map_fn, 1,
+ host_allocator);
+ map_call = fold_convert (TREE_TYPE (allocator), map_call);
+ gimplify_assign (orig_allocator, allocator, &pre_stmt_seq);
+ gimplify_assign (allocator, map_call, &pre_stmt_seq);
+
+ tree unmap_call = build_call_expr_loc (OMP_CLAUSE_LOCATION (c),
+ omp_allocator_unmap_fn,
1,
+ host_allocator);
+ gimplify_and_add (unmap_call, &post_stmt_seq);
+
+ call = build_call_expr_loc (OMP_CLAUSE_LOCATION (c),
+ omp_destroy_allocator_fn, 1,
+ host_allocator);
+ gimplify_and_add (call, &post_stmt_seq);
+ gimplify_assign (allocator, orig_allocator, &post_stmt_seq);
+ }
+
+ if (fini_seq)
+ {
+ gbind *bind = as_a<gbind *> (gimple_seq_first_stmt (body));
+ g = gimple_build_try (gimple_bind_body (bind),
+ fini_seq, GIMPLE_TRY_FINALLY);
+ gimple_seq_add_stmt (&init_seq, g);
+ gimple_bind_set_body (bind, init_seq);
+ body = bind;
+ }
+ }
}
else
gimplify_and_add (OMP_BODY (expr), &body);
@@ -18928,7 +19117,11 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq
*pre_p)
gcc_unreachable ();
}
+ if (pre_stmt_seq)
+ gimplify_seq_add_seq (pre_p, pre_stmt_seq);
gimplify_seq_add_stmt (pre_p, stmt);
+ if (post_stmt_seq)
+ gimplify_seq_add_seq (pre_p, post_stmt_seq);
*expr_p = NULL_TREE;
}
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index 010885f3f03..437b42308c1 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -93,6 +93,14 @@ DEF_GOMP_BUILTIN (BUILT_IN_OMP_SET_DEFAULT_DEVICE,
"omp_set_default_device",
BT_FN_INT, ATTR_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_INTEROP_INT, "omp_get_interop_int",
BT_FN_PTRMODE_PTR_INT_PTR, ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_OMP_INIT_ALLOCATOR, "omp_init_allocator",
+ BT_FN_PTRMODE_PTRMODE_INT_PTR, ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_OMP_DESTROY_ALLOCATOR, "omp_destroy_allocator",
+ BT_FN_VOID_PTRMODE, ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_OMP_ALLOCATOR_MAP, "GOMP_omp_allocator_map",
+ BT_FN_PTRMODE_PTRMODE, ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_OMP_ALLOCATOR_UNMAP,
"GOMP_omp_allocator_unmap",
+ BT_FN_VOID_PTRMODE, ATTR_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_ATOMIC_START, "GOMP_atomic_start",
BT_FN_VOID, ATTR_NOTHROW_LEAF_LIST)
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index d36756e33a5..a8837f5db72 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1173,6 +1173,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;
@@ -1768,6 +1798,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_FINALIZE:
case OMP_CLAUSE_TASK_REDUCTION:
case OMP_CLAUSE_ALLOCATE:
+ case OMP_CLAUSE_USES_ALLOCATORS:
break;
case OMP_CLAUSE_ALIGNED:
@@ -1994,6 +2025,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_INIT:
case OMP_CLAUSE_USE:
case OMP_CLAUSE_DESTROY:
+ case OMP_CLAUSE_USES_ALLOCATORS:
break;
case OMP_CLAUSE__CACHE_:
diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c
b/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c
new file mode 100644
index 00000000000..5a2e4a90e54
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c
@@ -0,0 +1,46 @@
+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);
+
+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" "" { 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" } */
+ a = 7;
+ }
+}
+
+void
+g (omp_allocator_handle_t my_alloc)
+{
+ /* The following defines a default-mem-space allocator with no extra traits.
*/
+ #pragma omp target uses_allocators(my_alloc)
+ {
+ int a;
+ #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)
+ a = 7;
+ }
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c
b/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c
new file mode 100644
index 00000000000..4dd1f13100a
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c
@@ -0,0 +1,33 @@
+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;
+
+void
+f ()
+{
+ omp_alloctrait_t trait[1] = {{1,1}};
+ omp_allocator_handle_t my_alloc;
+ #pragma omp target uses_allocators(traits(trait) : my_alloc) /* { dg-error
"traits array 'trait' must be of 'const omp_alloctrait_t \\\[\\\]' type" } */
+ ;
+}
+
+void
+g ()
+{
+ const omp_alloctrait_t trait[1] = {{1,1}};
+ omp_allocator_handle_t my_alloc;
+ #pragma omp target uses_allocators(traits(trait) : my_alloc)
+ ;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-3.c
b/gcc/testsuite/c-c++-common/gomp/uses_allocators-3.c
new file mode 100644
index 00000000000..e5fa1195e6b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-3.c
@@ -0,0 +1,58 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */
+
+#include <omp.h>
+
+omp_alloctrait_key_t k;
+omp_alloctrait_value_t v;
+
+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))
+ ;
+ #pragma omp target uses_allocators (foo (foo_traits), bar (foo_traits))
+ ;
+ #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\\(bar:
memspace\\(\\), traits\\(\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(foo:
memspace\\(\\), traits\\(foo_traits\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar:
memspace\\(\\), traits\\(foo_traits\\)\\) uses_allocators\\(foo:
memspace\\(\\), traits\\(foo_traits\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(foo:
memspace\\(.+\\), traits\\(\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar:
memspace\\(\\), traits\\(foo_traits\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar:
memspace\\(.+\\), traits\\(foo_traits\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar:
memspace\\(.+\\), traits\\(foo_traits\\)\\) uses_allocators\\(foo:
memspace\\(\\), traits\\(\\)\\)" "original" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\)
thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(\\), traits\\(\\)\\)
firstprivate\\(bar\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\)
thread_limit\\(0\\) uses_allocators\\(foo: memspace\\(\\),
traits\\(foo_traits\\)\\) firstprivate\\(foo\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\)
thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(\\),
traits\\(foo_traits\\)\\) uses_allocators\\(foo: memspace\\(\\),
traits\\(foo_traits\\)\\) firstprivate\\(bar\\) firstprivate\\(foo\\)" "gimple"
} } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\)
thread_limit\\(0\\) uses_allocators\\(foo: memspace\\(.+\\), traits\\(\\)\\)
firstprivate\\(foo\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\)
thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(\\),
traits\\(foo_traits\\)\\) firstprivate\\(bar\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\)
thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(.+\\),
traits\\(foo_traits\\)\\) firstprivate\\(bar\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\)
thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(.+\\),
traits\\(foo_traits\\)\\) uses_allocators\\(foo: memspace\\(\\),
traits\\(\\)\\) firstprivate\\(bar\\) firstprivate\\(foo\\)" "gimple" } } */
+
+/* { dg-final { scan-tree-dump-times "__builtin_omp_init_allocator" 9 "gimple"
} } */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_omp_allocator_map" 9
"gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_omp_allocator_unmap" 9
"gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_omp_destroy_allocator" 9
"gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-4.c
b/gcc/testsuite/c-c++-common/gomp/uses_allocators-4.c
new file mode 100644
index 00000000000..da289818e6d
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-4.c
@@ -0,0 +1,62 @@
+/* { dg-do compile } */
+
+#include <omp.h>
+
+omp_alloctrait_key_t k;
+omp_alloctrait_value_t v;
+
+int f (const omp_alloctrait_t arg_traits[], int n)
+{
+ omp_allocator_handle_t foo, bar;
+ const omp_alloctrait_t traits_array[] = { { omp_atk_pinned, omp_atv_true
},
+ { omp_atk_partition,
omp_atv_nearest } };
+ extern const omp_alloctrait_t ex_traits[2];
+ extern const omp_alloctrait_t ex2_traits[];
+#ifndef __cplusplus
+ const omp_alloctrait_t vla_traits[n] = {}; /* Not useful, but shouldn't
crash. */
+#else
+ const omp_alloctrait_t vla_traits[n] = { { omp_atk_pinned, omp_atv_true },
+ { omp_atk_partition, omp_atv_nearest
} };
+#endif
+
+ #pragma omp target uses_allocators (baz) /* { dg-error "'baz' undeclared
.first use in this function." "" { target c } } */
+ ; /* { dg-error "'baz' has not been
declared" "" { target c++ } .-1 } */
+ #pragma omp target uses_allocators (foo (xyz)) /* { dg-error "'xyz'
undeclared .first use in this function." "" { target c } } */
+ ; /* { dg-error "'xyz' has not
been declared" "" { target c++ } .-1 } */
+ #pragma omp target uses_allocators (foo (traits_array), baz (traits_array))
/* { dg-error "'baz' has not been declared" "" { target c++ } } */
+ ;
+ #pragma omp target uses_allocators (foo (arg_traits)) /* { dg-error "traits
array 'arg_traits' must be defined in same scope as the construct on which the
clause appears" } */
+ ; /* { dg-error "traits
array 'arg_traits' must be of 'const omp_alloctrait_t \\\[\\\]' type" "" {
target *-*-* } .-1 } */
+ #pragma omp target uses_allocators (foo (ex_traits)) /* { dg-error "traits
array 'ex_traits' must be defined in same scope as the construct on which the
clause appears" } */
+ ; /* { dg-error "traits
array must be initialized with constants" "" { target *-*-* } .-1 } */
+ #pragma omp target uses_allocators (foo (ex2_traits)) /* { dg-error "traits
array 'ex2_traits' must be defined in same scope as the construct on which the
clause appears" } */
+ ; /* { dg-error "traits
array 'ex2_traits' must be of 'const omp_alloctrait_t \\\[\\\]' type" "" {
target *-*-* } .-1 } */
+ #pragma omp target uses_allocators (foo (vla_traits)) /* { dg-error
"variable length traits arrays are not supported" "" { target c++ } } */
+ ;
+ #pragma omp target uses_allocators (memspace(omp_no_such_space) : foo) /* {
dg-error "'omp_no_such_space' undeclared .first use in this function." "" {
target c } } */
+ ; /* {
dg-error "'omp_no_such_space' was not declared in this scope" "" { target c++ }
.-1 } */
+ #pragma omp target uses_allocators (memspace(1) : foo) /* { dg-error
"memspace modifier must be constant enum of 'omp_memspace_handle_t' type" } */
+ ;
+ #pragma omp target uses_allocators (memspace(omp_no_such_space) : foo, bar)
/* { dg-error "'uses_allocators' clause only accepts a single allocator when
using modifiers" } */
+ ;
/* { dg-error "memspace modifier must be constant enum of
'omp_memspace_handle_t' type" "" { target c++ } .-1 } */
+ #pragma omp target uses_allocators (traits(xyz) : bar) /* { dg-error "'xyz'
was not declared in this scope" "" { target c++ } } */
+ ;
+ #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space),
traits(traits_array), memspace (omp_no_such_space) : bar) /* { dg-error
"duplicate 'memspace' modifier" "" { target c } } */
+ ;
/* { dg-error "expected
'\\\)' before 'memspace" "" { target c } .-1 } */
+
/* { dg-error "duplicate
'memspace' modifier" "" { target c++ } .-2 } */
+ #pragma omp target uses_allocators (traitz(traits_array),
memspace(omp_high_bw_mem_space) : bar) /* { dg-error "'traitz' undeclared
.first use in this function." "" { target c } } */
+ ;
/* { dg-error "'memspace' undeclared .first use in this
function." "" { target c } .-1 } */
+
/* { dg-error "'traitz' has not been declared" "" { target
c++ } .-2 } */
+
/* { dg-error "'memspace' has not been declared" "" {
target c++ } .-3 } */
+
/* { dg-error "expected '\\\)' before ':' token" "" {
target *-*-* } .-4 } */
+ #pragma omp target uses_allocators (omp_null_allocator)
+ ;
+ #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo,
bar) /* { dg-error "'uses_allocators' clause only accepts a single allocator
when using modifiers" } */
+ ;
+ #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space) :
foo(foo_traits)) /* { dg-error "'foo_traits' undeclared .first use in this
function.; did you mean 'vla_traits'." "" { target c } } */
+ ;
/* { dg-error "'foo_traits' has not been declared" "" { target c++ }
.-1 } */
+
/* { dg-error "legacy 'foo\\\(foo_traits\\\)' traits syntax not
allowed in 'uses_allocators' clause when using modifiers" "" { target *-*-* }
.-2 } */
+ return 0;
+}
+
+
diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-5.c
b/gcc/testsuite/c-c++-common/gomp/uses_allocators-5.c
new file mode 100644
index 00000000000..d3b9dbae2de
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-5.c
@@ -0,0 +1,39 @@
+/* { 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))
+ ;
+ #pragma omp target uses_allocators (memspace (mytraits),
omp_default_mem_alloc)
+ ;
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace:
memspace\\(\\), traits\\(\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(traits:
memspace\\(\\), traits\\(\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace:
memspace\\(\\), traits\\(\\)\\) uses_allocators\\(traits: memspace\\(\\),
traits\\(\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(traits:
memspace\\(\\), traits\\(mytraits\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace:
memspace\\(\\), traits\\(mytraits\\)\\)" "original" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\)
thread_limit\\(0\\) uses_allocators\\(memspace: memspace\\(\\), traits\\(\\)\\)
firstprivate\\(memspace\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\)
thread_limit\\(0\\) uses_allocators\\(traits: memspace\\(\\), traits\\(\\)\\)
firstprivate\\(traits\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\)
thread_limit\\(0\\) uses_allocators\\(memspace: memspace\\(\\), traits\\(\\)\\)
uses_allocators\\(traits: memspace\\(\\), traits\\(\\)\\)
firstprivate\\(traits\\) firstprivate\\(memspace\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\)
thread_limit\\(0\\) uses_allocators\\(traits: memspace\\(\\),
traits\\(mytraits\\)\\) firstprivate\\(traits\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\)
thread_limit\\(0\\) uses_allocators\\(memspace: memspace\\(\\),
traits\\(mytraits\\)\\) firstprivate\\(memspace\\)" "gimple" } } */
+
+/* { dg-final { scan-tree-dump-times "__builtin_omp_init_allocator" 6 "gimple"
} } */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_omp_allocator_map" 6
"gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_omp_allocator_unmap" 6
"gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_omp_destroy_allocator" 6
"gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-6.c
b/gcc/testsuite/c-c++-common/gomp/uses_allocators-6.c
new file mode 100644
index 00000000000..5942a0d6bbd
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/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
.*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\),
traits\\(trait\\)\\)" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target
.*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\),
traits\\(\\)\\)" 1 "gimple" } } */
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/testsuite/gfortran.dg/gomp/uses_allocators-1.f90
b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-1.f90
new file mode 100644
index 00000000000..8694cf5d03b
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-1.f90
@@ -0,0 +1,17 @@
+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))
+block; end block
+end
+end
diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_1.f90
b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-2.f90
similarity index 100%
rename from libgomp/testsuite/libgomp.fortran/uses_allocators_1.f90
rename to gcc/testsuite/gfortran.dg/gomp/uses_allocators-2.f90
diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90
b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-3.f90
similarity index 70%
rename from libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90
rename to gcc/testsuite/gfortran.dg/gomp/uses_allocators-3.f90
index 07327969775..bb984033413 100644
--- a/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/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/gcc/testsuite/gfortran.dg/gomp/uses_allocators-4.f90
b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-4.f90
new file mode 100644
index 00000000000..0458e338262
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-4.f90
@@ -0,0 +1,62 @@
+! { 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))
+ 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\\(bar: memspace\\(\\),
traits\\(\\)\\)" "original" } }
+! { dg-final { scan-tree-dump "pragma omp target" "original" } }
+! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar:
memspace\\(\\), traits\\(traits_array\\)\\) uses_allocators\\(foo:
memspace\\(\\), traits\\(traits_array\\)\\)" "original" } }
+! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar:
memspace\\(\\), traits\\(traits_array\\)\\)" "original" } }
+! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar:
memspace\\(omp_low_lat_mem_space\\), traits\\(\\)\\)" "original" } }
+! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar:
memspace\\(omp_high_bw_mem_space\\), traits\\(traits_array\\)\\)" "original" } }
+
+! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\)
thread_limit\\(0\\) allocate\\(allocator\\(bar\\):arr\\) uses_allocators\\(bar:
memspace\\(\\), traits\\(\\)\\) firstprivate\\(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\\(bar: memspace\\(\\),
traits\\(traits_array\\)\\) uses_allocators\\(foo: memspace\\(\\),
traits\\(traits_array\\)\\) firstprivate\\(foo\\) firstprivate\\(bar\\)"
"gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\)
thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(\\),
traits\\(traits_array\\)\\) firstprivate\\(bar\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\)
thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(omp_low_lat_mem_space\\),
traits\\(\\)\\) firstprivate\\(omp_low_lat_mem_space\\) firstprivate\\(bar\\)"
"gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\)
thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\),
traits\\(traits_array\\)\\) firstprivate\\(omp_high_bw_mem_space\\)
firstprivate\\(bar\\)" "gimple" } }
+
+! { dg-final { scan-tree-dump-times "__builtin_omp_init_allocator" 6 "gimple"
} }
+! { dg-final { scan-tree-dump-times "__builtin_omp_destroy_allocator" 6
"gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/uses_allocators-5.f90
b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-5.f90
new file mode 100644
index 00000000000..00f1dcb2763
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/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) ! { dg-error "Allocator
'omp_null_allocator' at .1. in USES_ALLOCATORS must either a variable or a
predefined 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/gcc/testsuite/gfortran.dg/gomp/uses_allocators-6.f90
b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-6.f90
new file mode 100644
index 00000000000..00f87109d2c
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/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/gcc/testsuite/gfortran.dg/gomp/uses_allocators-7.f90
b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-7.f90
new file mode 100644
index 00000000000..3799e3cec73
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-7.f90
@@ -0,0 +1,50 @@
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+ use iso_c_binding
+ use omp_lib
+ implicit none (type, external)
+ integer :: x, xbuf(10)
+ 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(omp_low_lat_mem_alloc) map(tofrom: x, xbuf)
defaultmap(none)
+ !$omp parallel allocate(allocator(omp_low_lat_mem_alloc), align(128): x,
xbuf) if(.false.) firstprivate(x, xbuf)
+ if (mod (TRANSFER (loc(x), iptr), 128) /= 0) &
+ stop 1
+ if (mod (TRANSFER (loc(xbuf), iptr), 128) /= 0) &
+ stop 2
+ !$omp end parallel
+ !$omp end target
+
+ 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 parallel allocate(allocator(my_alloc): x, xbuf) if(.false.)
firstprivate(x, xbuf)
+ if (mod (TRANSFER (loc(x), iptr), 128) /= 0) &
+ stop 3
+ if (mod (TRANSFER (loc(xbuf), iptr), 128) /= 0) &
+ stop 4
+ !$omp end parallel
+ !$omp end target
+
+ if (transfer(my_alloc, 0_omp_allocator_handle_kind) /= int(z'ABCD',
omp_allocator_handle_kind)) &
+ stop 5
+
+ ! The following creates an allocator with empty traits + default mem space.
+ !$omp target uses_allocators(my_alloc) map(tofrom: x, xbuf) defaultmap(none)
+ !$omp parallel allocate(allocator(my_alloc), align(128): x, xbuf)
if(.false.) firstprivate(x, xbuf)
+ if (mod (TRANSFER (loc(x), iptr), 128) /= 0) &
+ stop 6
+ if (mod (TRANSFER (loc(xbuf), iptr), 128) /= 0) &
+ stop 7
+ !$omp end parallel
+ !$omp end target
+
+ if (transfer(my_alloc, 0_omp_allocator_handle_kind) /= int(z'ABCD',
omp_allocator_handle_kind)) &
+ stop 8
+end
+
+! { dg-final { scan-tree-dump-times "#pragma omp target
.*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\),
traits\\(trait\\)\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target
.*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\),
traits\\(\\)\\)" 1 "gimple" } }
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 33a6a78f00d..6e55a079b3c 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -598,6 +598,8 @@ enum omp_clause_code {
/* OpenMP clause: nocontext (scalar-expression). */
OMP_CLAUSE_NOCONTEXT,
+ /* OpenMP clause: uses_allocators. */
+ OMP_CLAUSE_USES_ALLOCATORS,
};
#undef DEFTREESTRUCT
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index c19babadead..f5e4fbbfe0c 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -907,6 +907,20 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc,
dump_flags_t flags)
pp_right_paren (pp);
break;
+ case OMP_CLAUSE_USES_ALLOCATORS:
+ pp_string (pp, "uses_allocators(");
+ dump_generic_node (pp, OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (clause),
+ spc, flags, false);
+ pp_string (pp, ": memspace(");
+ dump_generic_node (pp, OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (clause),
+ spc, flags, false);
+ pp_string (pp, "), traits(");
+ dump_generic_node (pp, OMP_CLAUSE_USES_ALLOCATORS_TRAITS (clause),
+ spc, flags, false);
+ pp_right_paren (pp);
+ pp_right_paren (pp);
+ break;
+
case OMP_CLAUSE_AFFINITY:
pp_string (pp, "affinity(");
{
diff --git a/gcc/tree.cc b/gcc/tree.cc
index 4c8e31cfb12..cb48490d3fd 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -398,6 +398,7 @@ unsigned const char omp_clause_num_ops[] =
0, /* OMP_CLAUSE_NOHOST */
1, /* OMP_CLAUSE_NOVARIANTS */
1, /* OMP_CLAUSE_NOCONTEXT */
+ 3, /* OMP_CLAUSE_USES_ALLOCATORS */
};
const char * const omp_clause_code_name[] =
@@ -501,6 +502,7 @@ const char * const omp_clause_code_name[] =
"nohost",
"novariants",
"nocontext",
+ "uses_allocators",
};
/* Unless specific to OpenACC, we tend to internally maintain OpenMP-centric
diff --git a/gcc/tree.h b/gcc/tree.h
index 762228c336f..53c284b0887 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -2065,6 +2065,15 @@ class auto_suppress_location_wrappers
#define OMP_CLAUSE_ALLOCATE_COMBINED(NODE) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_ALLOCATE)->base.public_flag)
+#define OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE,
OMP_CLAUSE_USES_ALLOCATORS), 0)
+
+#define OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE,
OMP_CLAUSE_USES_ALLOCATORS), 1)
+
+#define OMP_CLAUSE_USES_ALLOCATORS_TRAITS(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE,
OMP_CLAUSE_USES_ALLOCATORS), 2)
+
#define OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR(NODE) \
OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_NUM_TEAMS), 0)
diff --git a/libgomp/allocator.c b/libgomp/allocator.c
index 4a683d90bba..bdb6f61b9a8 100644
--- a/libgomp/allocator.c
+++ b/libgomp/allocator.c
@@ -218,14 +218,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(OFFLOAD_PLUGINS) || 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;
@@ -1469,3 +1479,20 @@ fail:;
}
return NULL;
}
+
+#if !defined(LIBGOMP_OFFLOADED_ONLY)
+/* Called only from host-side at GOMP_omp_allocator_map, used to do memspace
+ validation using offload plugin. */
+attribute_hidden bool
+gomp_memspace_validate (struct gomp_device_descr *devicep, void *ptr)
+{
+ /* This is kept here to contain the definition of struct omp_allocator_data
+ within allocator.c. */
+ if (devicep->memspace_validate_func)
+ {
+ struct omp_allocator_data *data = (struct omp_allocator_data *) ptr;
+ return devicep->memspace_validate_func (data->memspace, data->access);
+ }
+ return true;
+}
+#endif
diff --git a/libgomp/config/nvptx/allocator.c b/libgomp/config/nvptx/allocator.c
index 7e9e343d2a9..cf751f84b6e 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;
@@ -62,12 +65,16 @@ nvptx_memspace_alloc (omp_memspace_handle_t memspace,
size_t size)
return __nvptx_lowlat_alloc (shared_pool, size);
}
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;
@@ -76,12 +83,16 @@ nvptx_memspace_calloc (omp_memspace_handle_t memspace,
size_t size)
return __nvptx_lowlat_calloc (shared_pool, size);
}
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;
@@ -90,6 +101,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);
}
@@ -97,6 +109,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;
@@ -105,22 +120,17 @@ nvptx_memspace_realloc (omp_memspace_handle_t memspace,
void *addr,
return __nvptx_lowlat_realloc (shared_pool, addr, oldsize, size);
}
else
+#endif
return realloc (addr, size);
}
static inline int
nvptx_memspace_validate (omp_memspace_handle_t memspace, unsigned access)
{
-#if __PTX_ISA_VERSION_MAJOR__ > 4 \
- || (__PTX_ISA_VERSION_MAJOR__ == 4 && __PTX_ISA_VERSION_MINOR >= 1)
/* 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);
-#else
- /* Low-latency memory is not available before PTX 4.1. */
- return (memspace != omp_low_lat_mem_space);
-#endif
}
#define MEMSPACE_ALLOC(MEMSPACE, SIZE, PIN) \
diff --git a/libgomp/configure b/libgomp/configure
index f522a6e5d27..8f79d698f56 100755
--- a/libgomp/configure
+++ b/libgomp/configure
@@ -15744,11 +15744,13 @@ if test x"$enable_offload_targets" != x; then
fi
done
fi
+if test x"$offload_plugins" != x; then
cat >>confdefs.h <<_ACEOF
#define OFFLOAD_PLUGINS "$offload_plugins"
_ACEOF
+fi
if test $PLUGIN_NVPTX = 1; then
PLUGIN_NVPTX_TRUE=
PLUGIN_NVPTX_FALSE='#'
@@ -17100,6 +17102,17 @@ $as_echo "#define HAVE_SYNC_BUILTINS 1" >>confdefs.h
fi
+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 1730c62c74c..be4665d5391 100644
--- a/libgomp/configure.ac
+++ b/libgomp/configure.ac
@@ -365,6 +365,17 @@ CFLAGS="$save_CFLAGS $XCFLAGS"
# had a chance to set XCFLAGS.
LIBGOMP_CHECK_SYNC_BUILTINS
+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 f2baed9bad9..46773a6e498 100644
--- a/libgomp/libgomp-plugin.h
+++ b/libgomp/libgomp-plugin.h
@@ -184,6 +184,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 *);
extern void GOMP_OFFLOAD_run (int, void *, void *, void **);
extern void GOMP_OFFLOAD_async_run (int, void *, void *, void **, void *);
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 3d406be175e..519768b5fee 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -614,6 +614,7 @@ extern bool gomp_display_affinity_var;
extern char *gomp_affinity_format_var;
extern size_t gomp_affinity_format_len;
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;
@@ -1041,6 +1042,10 @@ extern void gomp_display_affinity_thread
(gomp_thread_handle,
struct gomp_team_state *,
unsigned int) __attribute__((cold));
+/* allocator.c */
+
+extern bool gomp_memspace_validate (struct gomp_device_descr *, void *);
+
/* env.c */
extern struct gomp_icv_list *gomp_get_initial_icv_item (int dev_num);
@@ -1429,6 +1434,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;
__typeof (GOMP_OFFLOAD_can_run) *can_run_func;
__typeof (GOMP_OFFLOAD_run) *run_func;
__typeof (GOMP_OFFLOAD_async_run) *async_run_func;
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 67e08a37116..570fe6c693a 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -457,6 +457,8 @@ GOMP_6.0.1 {
global:
omp_target_memset;
omp_target_memset_async;
+ GOMP_omp_allocator_map;
+ GOMP_omp_allocator_unmap;
} GOMP_6.0;
OACC_2.0 {
diff --git a/libgomp/oacc-host.c b/libgomp/oacc-host.c
index cd09165cfe3..e0fc53e2712 100644
--- a/libgomp/oacc-host.c
+++ b/libgomp/oacc-host.c
@@ -286,6 +286,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,
.mem_map = { NULL },
diff --git a/libgomp/plugin/configfrag.ac b/libgomp/plugin/configfrag.ac
index e7a69975baa..dc3d4cb05c7 100644
--- a/libgomp/plugin/configfrag.ac
+++ b/libgomp/plugin/configfrag.ac
@@ -123,7 +123,9 @@ if test x"$enable_offload_targets" != x; then
fi
done
fi
-AC_DEFINE_UNQUOTED(OFFLOAD_PLUGINS, "$offload_plugins",
- [Define to offload plugins, separated by commas.])
+if test x"$offload_plugins" != x; then
+ AC_DEFINE_UNQUOTED(OFFLOAD_PLUGINS, "$offload_plugins",
+ [Define to offload plugins, separated by commas.])
+fi
AM_CONDITIONAL([PLUGIN_NVPTX], [test $PLUGIN_NVPTX = 1])
AM_CONDITIONAL([PLUGIN_GCN], [test $PLUGIN_GCN = 1])
diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c
index cd5a19b0355..c5a6695efff 100644
--- a/libgomp/plugin/plugin-gcn.c
+++ b/libgomp/plugin/plugin-gcn.c
@@ -4496,6 +4496,14 @@ unlock:
return retval;
}
+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);
+}
static bool
init_hip_runtime_functions (void)
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 5ad66688e7e..fd3afd1b3ad 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -2397,6 +2397,15 @@ GOMP_OFFLOAD_memset (int ord, void *ptr, int val, size_t
count)
return true;
}
+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);
+}
+
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 002a144b4ab..a6c7dffba30 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -5916,6 +5916,62 @@ omp_get_device_from_uid (const char *uid)
ialias (omp_get_uid_from_device)
ialias (omp_get_device_from_uid)
+omp_allocator_handle_t
+GOMP_omp_allocator_map (omp_allocator_handle_t host_handle)
+{
+ if (!host_handle)
+ return host_handle;
+ struct gomp_device_descr *devicep = resolve_device (-1, true);
+ if (devicep == NULL
+ || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
+ || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return host_handle;
+
+ if (!gomp_memspace_validate (devicep, (void *) host_handle))
+ {
+ free ((void *) host_handle);
+ return omp_null_allocator;
+ }
+
+ unsigned short kind = GOMP_MAP_TO;
+ size_t size = gomp_omp_allocator_data_size;
+
+ struct omp_allocator_data *ptr
+ = (struct omp_allocator_data *) host_handle;
+
+ gomp_map_vars (devicep, 1, (void **) &ptr, NULL, &size, &kind, true, NULL,
+ GOMP_MAP_VARS_ENTER_DATA);
+ struct splay_tree_key_s node;
+ node.host_start = (uintptr_t) ptr;
+ node.host_end = node.host_start + size;
+ gomp_mutex_lock (&devicep->lock);
+ splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &node);
+ gomp_mutex_unlock (&devicep->lock);
+ if (!n)
+ gomp_fatal ("Mapped allocator not found on device");
+ return (omp_allocator_handle_t) n->tgt->tgt_start;
+}
+
+void
+GOMP_omp_allocator_unmap (omp_allocator_handle_t host_handle)
+{
+ if (!host_handle)
+ return;
+ struct gomp_device_descr *devicep = resolve_device (-1, true);
+ if (devicep == NULL
+ || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
+ || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return;
+
+ unsigned short kind = GOMP_MAP_DELETE;
+ size_t size = gomp_omp_allocator_data_size;
+
+ struct omp_allocator_data *ptr
+ = (struct omp_allocator_data *) host_handle;
+
+ gomp_exit_data (devicep, 1, (void **) &ptr, &size, &kind, NULL);
+}
+
#ifdef PLUGIN_SUPPORT
/* This function tries to load a plugin for DEVICE. Name of plugin is passed
@@ -5973,6 +6029,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/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-7.f90
b/libgomp/testsuite/libgomp.fortran/uses_allocators-7.f90
new file mode 100644
index 00000000000..3799e3cec73
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/uses_allocators-7.f90
@@ -0,0 +1,50 @@
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+ use iso_c_binding
+ use omp_lib
+ implicit none (type, external)
+ integer :: x, xbuf(10)
+ 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(omp_low_lat_mem_alloc) map(tofrom: x, xbuf)
defaultmap(none)
+ !$omp parallel allocate(allocator(omp_low_lat_mem_alloc), align(128): x,
xbuf) if(.false.) firstprivate(x, xbuf)
+ if (mod (TRANSFER (loc(x), iptr), 128) /= 0) &
+ stop 1
+ if (mod (TRANSFER (loc(xbuf), iptr), 128) /= 0) &
+ stop 2
+ !$omp end parallel
+ !$omp end target
+
+ 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 parallel allocate(allocator(my_alloc): x, xbuf) if(.false.)
firstprivate(x, xbuf)
+ if (mod (TRANSFER (loc(x), iptr), 128) /= 0) &
+ stop 3
+ if (mod (TRANSFER (loc(xbuf), iptr), 128) /= 0) &
+ stop 4
+ !$omp end parallel
+ !$omp end target
+
+ if (transfer(my_alloc, 0_omp_allocator_handle_kind) /= int(z'ABCD',
omp_allocator_handle_kind)) &
+ stop 5
+
+ ! The following creates an allocator with empty traits + default mem space.
+ !$omp target uses_allocators(my_alloc) map(tofrom: x, xbuf) defaultmap(none)
+ !$omp parallel allocate(allocator(my_alloc), align(128): x, xbuf)
if(.false.) firstprivate(x, xbuf)
+ if (mod (TRANSFER (loc(x), iptr), 128) /= 0) &
+ stop 6
+ if (mod (TRANSFER (loc(xbuf), iptr), 128) /= 0) &
+ stop 7
+ !$omp end parallel
+ !$omp end target
+
+ if (transfer(my_alloc, 0_omp_allocator_handle_kind) /= int(z'ABCD',
omp_allocator_handle_kind)) &
+ stop 8
+end
+
+! { dg-final { scan-tree-dump-times "#pragma omp target
.*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\),
traits\\(trait\\)\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target
.*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\),
traits\\(\\)\\)" 1 "gimple" } }
