* [patch] OpenMP: Add uses_allocators support
@ 2023-11-20 10:42 Tobias Burnus
2023-12-08 13:15 ` Jakub Jelinek
0 siblings, 1 reply; 2+ messages in thread
From: Tobias Burnus @ 2023-11-20 10:42 UTC (permalink / raw)
To: gcc-patches, fortran
[-- Attachment #1: Type: text/plain, Size: 716 bytes --]
This adds middle end support for uses_allocators, wires Fortran to use
it and add C/C++ parsing support.
This is based on Chung-Lin's patch at
https://gcc.gnu.org/pipermail/gcc-patches/2022-June/596412.html
There is a known firstprivate/private issue for C/C++, see:
https://gcc.gnu.org/PR110347 but that one is unrelated.
Tested on x86-64-gnu-linux w/o offloading; with offloading still running.
Comments, suggestions?
Tobias
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
[-- Attachment #2: omp-uses-allocators.diff --]
[-- Type: text/x-patch, Size: 78473 bytes --]
OpenMP: Add uses_allocators support
2023-11-19 Tobias Burnus <tobias@codesourcery.com>
Chung-Lin Tang <cltang@codesourcery.com>
gcc/ChangeLog:
* builtin-types.def (BT_FN_VOID_PTRMODE):
(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,
BUILT_IN_OMP_DESTROY_ALLOCATOR): Add.
* omp-low.cc (scan_sharing_clauses):
* 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.
(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.
(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): Set PURPOSE
when building constructor for get_initialized_tmp_var.
* trans-openmp.cc (gfc_trans_omp_clauses): Handle uses_allocators.
* types.def (BT_FN_VOID_PTRMODE, BT_FN_PTRMODE_PTRMODE_INT_PTR): Add.
libgomp/ChangeLog:
* testsuite/libgomp.c++/c++.exp (check_effective_target_c,
check_effective_target_c++): Add.
* testsuite/libgomp.c/c.exp (check_effective_target_c,
check_effective_target_c++): Add.
* testsuite/libgomp.fortran/uses_allocators_2.f90: Remove 'sorry'.
* testsuite/libgomp.c-c++-common/uses_allocators-1.c: New test.
* testsuite/libgomp.c-c++-common/uses_allocators-2.c: New test.
* testsuite/libgomp.c-c++-common/uses_allocators-3.c: New test.
* testsuite/libgomp.c-c++-common/uses_allocators-4.c: New test.
* testsuite/libgomp.fortran/uses_allocators_3.f90: New test.
* testsuite/libgomp.fortran/uses_allocators_4.f90: New test.
* testsuite/libgomp.fortran/uses_allocators_5.f90: New test.
* testsuite/libgomp.fortran/uses_allocators_6.f90: New test.
gcc/testsuite/ChangeLog:
* gfortran.dg/gomp/allocate-1.f90: Add uses_allocators.
* gfortran.dg/gomp/scope-6.f90: Update dg-scan-tree-dump.
* c-c++-common/gomp/uses_allocators-1.c: New test.
* c-c++-common/gomp/uses_allocators-2.c: New test.
* gfortran.dg/gomp/uses_allocators-1.f90: New test.
Co-authored-by: Chung-Lin Tang <cltang@codesourcery.com>
gcc/builtin-types.def | 3 +
gcc/c-family/c-omp.cc | 1 +
gcc/c-family/c-pragma.h | 1 +
gcc/c/c-parser.cc | 216 ++++++++++++++++++-
gcc/c/c-typeck.cc | 105 +++++++++
gcc/cp/parser.cc | 237 ++++++++++++++++++++-
gcc/cp/semantics.cc | 95 +++++++++
gcc/fortran/trans-array.cc | 5 +-
gcc/fortran/trans-openmp.cc | 42 +++-
gcc/fortran/types.def | 3 +
gcc/gimplify.cc | 183 +++++++++++++++-
gcc/omp-builtins.def | 4 +
gcc/omp-low.cc | 32 +++
.../c-c++-common/gomp/uses_allocators-1.c | 46 ++++
.../c-c++-common/gomp/uses_allocators-2.c | 33 +++
gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 | 7 +-
gcc/testsuite/gfortran.dg/gomp/scope-6.f90 | 2 +-
.../gfortran.dg/gomp/uses_allocators-1.f90 | 23 ++
gcc/tree-core.h | 3 +
gcc/tree-pretty-print.cc | 14 ++
gcc/tree.cc | 2 +
gcc/tree.h | 9 +
libgomp/testsuite/libgomp.c++/c++.exp | 9 +
.../libgomp.c-c++-common/uses_allocators-1.c | 53 +++++
.../libgomp.c-c++-common/uses_allocators-2.c | 39 ++++
.../libgomp.c-c++-common/uses_allocators-3.c | 37 ++++
.../libgomp.c-c++-common/uses_allocators-4.c | 53 +++++
libgomp/testsuite/libgomp.c/c.exp | 8 +
.../libgomp.fortran/uses_allocators_2.f90 | 22 +-
.../libgomp.fortran/uses_allocators_3.f90 | 62 ++++++
.../libgomp.fortran/uses_allocators_4.f90 | 54 +++++
.../libgomp.fortran/uses_allocators_5.f90 | 14 ++
.../libgomp.fortran/uses_allocators_6.f90 | 50 +++++
33 files changed, 1435 insertions(+), 32 deletions(-)
diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index 43381bc8949..da9719272ec 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -385,6 +385,7 @@ DEF_FUNCTION_TYPE_1 (BT_FN_DFLOAT32_DFLOAT32, BT_DFLOAT32, BT_DFLOAT32)
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_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)
@@ -829,6 +830,8 @@ DEF_FUNCTION_TYPE_3 (BT_FN_PTR_SIZE_SIZE_PTRMODE,
BT_PTR, BT_SIZE, BT_SIZE, BT_PTRMODE)
DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_UINT8_PTRMODE, BT_VOID, BT_PTR, BT_UINT8,
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_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 95b6c1e623f..82422fc8c3e 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -2026,6 +2026,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 98177913053..7b1d9dc7232 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -161,6 +161,7 @@ enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_UNTIED,
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 703f9570dbc..db69a418f59 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -14722,6 +14722,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))
@@ -17568,6 +17570,213 @@ 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 t = NULL_TREE, nl = list;
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
+
+ tree memspace_expr = NULL_TREE;
+ tree traits_var = NULL_TREE;
+
+ struct item_tok
+ {
+ location_t loc;
+ tree id;
+ item_tok (void) : loc (UNKNOWN_LOCATION), id (NULL_TREE) {}
+ };
+ struct item { item_tok name, arg; };
+ auto_vec<item> *modifiers = NULL, *allocators = NULL;
+ auto_vec<item> *cur_list = new auto_vec<item> (4);
+
+ while (true)
+ {
+ item it;
+
+ if (c_parser_next_token_is (parser, CPP_NAME))
+ {
+ c_token *tok = c_parser_peek_token (parser);
+ it.name.id = tok->value;
+ it.name.loc = tok->location;
+ c_parser_consume_token (parser);
+
+ if (c_parser_next_token_is (parser, CPP_OPEN_PAREN))
+ {
+ matching_parens parens2;
+ parens2.consume_open (parser);
+
+ if (c_parser_next_token_is (parser, CPP_NAME))
+ {
+ tok = c_parser_peek_token (parser);
+ it.arg.id = tok->value;
+ it.arg.loc = tok->location;
+ c_parser_consume_token (parser);
+ }
+ else
+ {
+ c_parser_error (parser, "expected identifier");
+ parens2.skip_until_found_close (parser);
+ goto end;
+ }
+ parens2.skip_until_found_close (parser);
+ }
+ }
+
+ cur_list->safe_push (it);
+
+ if (c_parser_next_token_is (parser, CPP_COMMA))
+ c_parser_consume_token (parser);
+ else if (c_parser_next_token_is (parser, CPP_COLON))
+ {
+ if (modifiers)
+ {
+ c_parser_error (parser, "expected %<)%>");
+ goto end;
+ }
+ else
+ {
+ c_parser_consume_token (parser);
+ modifiers = cur_list;
+ cur_list = new auto_vec<item> (4);
+ }
+ }
+ else if (c_parser_next_token_is (parser, CPP_CLOSE_PAREN))
+ {
+ gcc_assert (allocators == NULL);
+ allocators = cur_list;
+ cur_list = NULL;
+ break;
+ }
+ else
+ {
+ c_parser_error (parser, "expected %<)%>");
+ goto end;
+ }
+ }
+
+ if (modifiers)
+ for (unsigned i = 0; i < modifiers->length (); i++)
+ {
+ item& it = (*modifiers)[i];
+ const char *p = IDENTIFIER_POINTER (it.name.id);
+ int strcmp_traits = 1, strcmp_memspace = 1;
+
+ if ((strcmp_traits = strcmp ("traits", p)) == 0
+ || (strcmp_memspace = strcmp ("memspace", p)) == 0)
+ {
+ if ((strcmp_traits == 0 && traits_var != NULL_TREE)
+ || (strcmp_memspace == 0 && memspace_expr != NULL_TREE))
+ {
+ error_at (it.name.loc, "duplicate %qs modifier", p);
+ goto end;
+ }
+ t = lookup_name (it.arg.id);
+ if (t == NULL_TREE)
+ {
+ undeclared_variable (it.arg.loc, it.arg.id);
+ t = error_mark_node;
+ }
+ else if (strcmp_memspace == 0)
+ memspace_expr = t;
+ else if (strcmp_traits == 0)
+ traits_var = t;
+ else
+ gcc_unreachable ();
+ }
+ else
+ {
+ error_at (it.name.loc, "unknown modifier %qE", it.name.id);
+ goto end;
+ }
+ }
+
+ if (allocators)
+ {
+ if (modifiers)
+ {
+ if (allocators->length () > 1)
+ {
+ error_at ((*allocators)[1].name.loc,
+ "%<uses_allocators%> clause only accepts a single "
+ "allocator when using modifiers");
+ goto end;
+ }
+ else if ((*allocators)[0].arg.id)
+ {
+ error_at ((*allocators)[0].arg.loc,
+ "legacy %<%E(%E)%> traits syntax not allowed in "
+ "%<uses_allocators%> clause when using modifiers",
+ (*allocators)[0].name.id, (*allocators)[0].arg.id);
+ goto end;
+ }
+ }
+
+ for (unsigned i = 0; i < allocators->length (); i++)
+ {
+ item& it = (*allocators)[i];
+ t = lookup_name (it.name.id);
+ if (t == NULL_TREE)
+ {
+ undeclared_variable (it.name.loc, it.name.id);
+ goto end;
+ }
+ else if (t != error_mark_node)
+ {
+ tree t2 = NULL_TREE;
+ if (it.arg.id)
+ {
+ t2 = lookup_name (it.arg.id);
+ if (t2 == NULL_TREE)
+ {
+ undeclared_variable (it.arg.loc, it.arg.id);
+ goto end;
+ }
+ }
+ else
+ t2 = traits_var;
+
+ 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) = t2;
+ OMP_CLAUSE_CHAIN (c) = nl;
+ nl = c;
+ }
+ }
+ }
+ end:
+ if (cur_list)
+ delete cur_list;
+ if (modifiers)
+ delete modifiers;
+ if (allocators)
+ delete allocators;
+ parens.skip_until_found_close (parser);
+ return nl;
+}
+
/* OpenMP 4.0:
linear ( variable-list )
linear ( variable-list : expression )
@@ -19237,6 +19446,10 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_omp_clause_allocate (parser, clauses);
c_name = "allocate";
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_LINEAR:
clauses = c_parser_omp_clause_linear (parser, clauses);
c_name = "linear";
@@ -23648,7 +23861,8 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
| (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
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 1dbb4471a88..e17414a5245 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -15136,6 +15136,111 @@ 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 ((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 (strcmp (IDENTIFIER_POINTER (DECL_NAME (t)),
+ "omp_null_allocator") == 0)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<omp_null_allocator%> cannot be used in "
+ "%<uses_allocators%> clause");
+ break;
+ }
+
+ 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 != NULL_TREE
+ && (TREE_CODE (t) != CONST_DECL
+ || 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 != NULL_TREE)
+ {
+ bool type_err = false;
+
+ if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE
+ || DECL_SIZE (t) == NULL_TREE)
+ 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 (TREE_CODE (t) != ERROR_MARK)
+ 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 "
+ "of constant values");
+
+ remove = true;
+ }
+ }
+ }
+ if (remove)
+ break;
+ pc = &OMP_CLAUSE_CHAIN (c);
+ continue;
case OMP_CLAUSE_DEPEND:
case OMP_CLAUSE_AFFINITY:
t = OMP_CLAUSE_DECL (c);
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index f556b8f3c01..15216a08b54 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -37657,6 +37657,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))
@@ -39983,6 +39985,234 @@ 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 t = NULL_TREE, nl = list;
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
+
+ tree memspace_expr = NULL_TREE;
+ tree traits_var = NULL_TREE;
+
+ struct item_tok
+ {
+ location_t loc;
+ tree id;
+ item_tok (void) : loc (UNKNOWN_LOCATION), id (NULL_TREE) {}
+ };
+ struct item { item_tok name, arg; };
+ auto_vec<item> *modifiers = NULL, *allocators = NULL;
+ auto_vec<item> *cur_list = new auto_vec<item> (4);
+
+ while (true)
+ {
+ item it;
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ {
+ cp_token *tok = cp_lexer_peek_token (parser->lexer);
+ it.name.id = tok->u.value;
+ it.name.loc = tok->location;
+ cp_lexer_consume_token (parser->lexer);
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_PAREN))
+ {
+ matching_parens parens2;
+ parens2.consume_open (parser);
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ {
+ tok = cp_lexer_peek_token (parser->lexer);
+ it.arg.id = tok->u.value;
+ it.arg.loc = tok->location;
+ cp_lexer_consume_token (parser->lexer);
+ }
+ else
+ {
+ cp_parser_error (parser, "expected identifier");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ goto end;
+ }
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/false,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ }
+ }
+
+ cur_list->safe_push (it);
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
+ cp_lexer_consume_token (parser->lexer);
+ else if (cp_lexer_next_token_is (parser->lexer, CPP_COLON))
+ {
+ if (modifiers)
+ {
+ cp_parser_error (parser, "expected %<)%>");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ goto end;
+ }
+ else
+ {
+ cp_lexer_consume_token (parser->lexer);
+ modifiers = cur_list;
+ cur_list = new auto_vec<item> (4);
+ }
+ }
+ else if (cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_PAREN))
+ {
+ gcc_assert (allocators == NULL);
+ allocators = cur_list;
+ cur_list = NULL;
+ break;
+ }
+ else
+ {
+ cp_parser_error (parser, "expected %<)%>");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ goto end;
+ }
+ }
+
+ if (modifiers)
+ for (unsigned i = 0; i < modifiers->length (); i++)
+ {
+ item& it = (*modifiers)[i];
+ const char *p = IDENTIFIER_POINTER (it.name.id);
+ int strcmp_traits = 1, strcmp_memspace = 1;
+
+ if ((strcmp_traits = strcmp ("traits", p)) == 0
+ || (strcmp_memspace = strcmp ("memspace", p)) == 0)
+ {
+ if ((strcmp_traits == 0 && traits_var != NULL_TREE)
+ || (strcmp_memspace == 0 && memspace_expr != NULL_TREE))
+ {
+ error_at (it.name.loc, "duplicate %qs modifier", p);
+ goto end;
+ }
+ t = cp_parser_lookup_name_simple (parser, it.arg.id, it.arg.loc);
+ if (t == error_mark_node)
+ {
+ cp_parser_name_lookup_error (parser, it.arg.id, t, NLE_NULL,
+ it.arg.loc);
+ }
+ else if (strcmp_memspace == 0)
+ memspace_expr = t;
+ else if (strcmp_traits == 0)
+ traits_var = t;
+ else
+ gcc_unreachable ();
+ }
+ else
+ {
+ error_at (it.name.loc, "unknown modifier %qE", it.name.id);
+ goto end;
+ }
+ }
+
+ if (allocators)
+ {
+ if (modifiers)
+ {
+ if (allocators->length () > 1)
+ {
+ error_at ((*allocators)[1].name.loc,
+ "%<uses_allocators%> clause only accepts a single "
+ "allocator when using modifiers");
+ goto end;
+ }
+ else if ((*allocators)[0].arg.id)
+ {
+ error_at ((*allocators)[0].arg.loc,
+ "legacy %<%E(%E)%> traits syntax not allowed in "
+ "%<uses_allocators%> clause when using modifiers",
+ (*allocators)[0].name.id, (*allocators)[0].arg.id);
+ goto end;
+ }
+ }
+
+ for (unsigned i = 0; i < allocators->length (); i++)
+ {
+ item& it = (*allocators)[i];
+ t = cp_parser_lookup_name_simple (parser, it.name.id, it.name.loc);
+ if (t == error_mark_node)
+ {
+ cp_parser_name_lookup_error (parser, it.name.id, t, NLE_NULL,
+ it.name.loc);
+ goto end;
+ }
+ else if (t != error_mark_node)
+ {
+ tree t2 = NULL_TREE;
+ if (it.arg.id)
+ {
+ t2 = cp_parser_lookup_name_simple (parser, it.arg.id,
+ it.arg.loc);
+ if (t2 == error_mark_node)
+ {
+ cp_parser_name_lookup_error (parser, it.arg.id, t2,
+ NLE_NULL, it.arg.loc);
+ goto end;
+ }
+ }
+ else
+ t2 = traits_var;
+
+ 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) = t2;
+ OMP_CLAUSE_CHAIN (c) = nl;
+ nl = c;
+ }
+ }
+ }
+ end:
+ if (cur_list)
+ delete cur_list;
+ if (modifiers)
+ delete modifiers;
+ if (allocators)
+ delete allocators;
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/false,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return nl;
+}
+
/* OpenMP 2.5:
lastprivate ( variable-list )
@@ -41856,6 +42086,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;
@@ -46427,7 +46661,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 8090c71809f..47796ab68c1 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -8029,6 +8029,101 @@ 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 (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 (strcmp (IDENTIFIER_POINTER (DECL_NAME (t)),
+ "omp_null_allocator") == 0)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<omp_null_allocator%> cannot be used in "
+ "%<uses_allocators%> clause");
+ break;
+ }
+
+ 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 != NULL_TREE
+ && (TREE_CODE (t) != CONST_DECL
+ || 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 != NULL_TREE)
+ {
+ bool type_err = false;
+
+ if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE
+ || DECL_SIZE (t) == NULL_TREE)
+ 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
+ {
+ tree cst_val = decl_constant_value (t);
+ if (cst_val == t)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "traits array must be "
+ "of constant values");
+
+ remove = true;
+ }
+ }
+ }
+ if (remove)
+ break;
+ pc = &OMP_CLAUSE_CHAIN (c);
+ continue;
case OMP_CLAUSE_DEPEND:
case OMP_CLAUSE_AFFINITY:
t = OMP_CLAUSE_DECL (c);
diff --git a/gcc/fortran/trans-array.cc b/gcc/fortran/trans-array.cc
index bbb81f40aa9..3329f8f3729 100644
--- a/gcc/fortran/trans-array.cc
+++ b/gcc/fortran/trans-array.cc
@@ -6501,10 +6501,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)
{
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 82bbc41b388..9905c882d29 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -2758,7 +2758,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_;
@@ -3938,13 +3943,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 7a465c89c5f..a9374d68d78 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -81,6 +81,7 @@ 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)
@@ -154,6 +155,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_VOID_OMPFN_PTR_UINT_UINT,
BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT, BT_UINT)
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index d52d71b9b6b..08bd37e2e22 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -1381,18 +1381,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),
@@ -12240,6 +12268,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_PRIVATE);
+
+ 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:
@@ -12375,6 +12418,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)
{
@@ -13411,6 +13497,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
case OMP_CLAUSE_FINALIZE:
case OMP_CLAUSE_INCLUSIVE:
case OMP_CLAUSE_EXCLUSIVE:
+ case OMP_CLAUSE_USES_ALLOCATORS:
break;
case OMP_CLAUSE_NOHOST:
@@ -15777,6 +15864,92 @@ 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);
+ tree t = DECL_INITIAL (traits);
+ gcc_assert (TREE_CODE (t) == CONSTRUCTOR);
+ ntraits = build_int_cst (integer_type_node,
+ CONSTRUCTOR_NELTS (t));
+ t = get_initialized_tmp_var (t, &init_seq, NULL);
+ traits_var = build_fold_addr_expr_loc (loc, t);
+ }
+ else
+ {
+ location_t loc = OMP_CLAUSE_LOCATION (c);
+ gcc_assert (TREE_CODE (TREE_TYPE (traits)) == ARRAY_TYPE);
+ tree t = TYPE_DOMAIN (TREE_TYPE (traits));
+ tree min = TYPE_MIN_VALUE (t);
+ tree max = TYPE_MAX_VALUE (t);
+ gcc_assert (TREE_CODE (min) == INTEGER_CST
+ && TREE_CODE (max) == INTEGER_CST);
+ t = fold_build2_loc (loc, MINUS_EXPR, TREE_TYPE (min),
+ max, min);
+ t = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (min),
+ t, build_int_cst (TREE_TYPE (min), 1));
+ ntraits = t;
+ 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 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 (allocator, call, &init_seq);
+
+ call = build_call_expr_loc (OMP_CLAUSE_LOCATION (c),
+ omp_destroy_allocator_fn, 1,
+ allocator);
+ gimplify_and_add (call, &fini_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);
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index ed78d49d205..c07f8533c0a 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -76,6 +76,10 @@ DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_TEAM_NUM, "omp_get_team_num",
BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_NUM_TEAMS, "omp_get_num_teams",
BT_FN_INT, ATTR_CONST_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_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 161bcfeec05..a53a620fc1f 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1166,6 +1166,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;
@@ -1756,6 +1786,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:
@@ -1974,6 +2005,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_FINALIZE:
case OMP_CLAUSE_FILTER:
case OMP_CLAUSE__CONDTEMP_:
+ 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/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..b20515b3691
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-1.f90
@@ -0,0 +1,23 @@
+use iso_c_binding
+implicit none
+ integer, parameter :: omp_allocator_handle_kind = c_intptr_t
+ integer, parameter :: omp_alloctrait_key_kind = c_int
+ integer, parameter :: omp_alloctrait_val_kind = c_intptr_t
+ type omp_alloctrait
+ integer (kind=omp_alloctrait_key_kind) key
+ integer (kind=omp_alloctrait_val_kind) value
+ end type omp_alloctrait
+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/gcc/tree-core.h b/gcc/tree-core.h
index 65e51b939a2..c76e7d6116f 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -533,6 +533,9 @@ enum omp_clause_code {
/* OpenACC clause: nohost. */
OMP_CLAUSE_NOHOST,
+
+ /* 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 1fadd752d05..8cf479d8508 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -800,6 +800,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 a3d907a4ac0..46f41cb7167 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -328,6 +328,7 @@ unsigned const char omp_clause_num_ops[] =
0, /* OMP_CLAUSE_IF_PRESENT */
0, /* OMP_CLAUSE_FINALIZE */
0, /* OMP_CLAUSE_NOHOST */
+ 3, /* OMP_CLAUSE_USES_ALLOCATORS */
};
const char * const omp_clause_code_name[] =
@@ -421,6 +422,7 @@ const char * const omp_clause_code_name[] =
"if_present",
"finalize",
"nohost",
+ "uses_allocators",
};
/* Unless specific to OpenACC, we tend to internally maintain OpenMP-centric
diff --git a/gcc/tree.h b/gcc/tree.h
index 086b55f0375..88befefe110 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1958,6 +1958,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/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++-common/uses_allocators-1.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-1.c
new file mode 100644
index 00000000000..eaca686d037
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-1.c
@@ -0,0 +1,53 @@
+/* { 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 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\\(omp_high_bw_mem_space\\), 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\\(omp_high_bw_mem_space\\), traits\\(foo_traits\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), 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\\(\\)\\) private\\(bar\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(foo: memspace\\(\\), traits\\(foo_traits\\)\\) private\\(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\\)\\) private\\(bar\\) private\\(foo\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(foo: memspace\\(omp_high_bw_mem_space\\), traits\\(\\)\\) private\\(foo\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(\\), traits\\(foo_traits\\)\\) private\\(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\\(foo_traits\\)\\) private\\(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\\(foo_traits\\)\\) uses_allocators\\(foo: memspace\\(\\), traits\\(\\)\\) private\\(bar\\) private\\(foo\\)" "gimple" } } */
+
+/* { dg-final { scan-tree-dump-times "__builtin_omp_init_allocator" 9 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_omp_destroy_allocator" 9 "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-2.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-2.c
new file mode 100644
index 00000000000..f350c0a409e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-2.c
@@ -0,0 +1,39 @@
+/* { dg-do compile } */
+
+#include <omp.h>
+
+omp_alloctrait_key_t k;
+omp_alloctrait_value_t v;
+
+int main (void)
+{
+ omp_allocator_handle_t foo, bar;
+ const omp_alloctrait_t traits_array[] = { { omp_atk_pinned, omp_atv_true },
+ { omp_atk_partition, omp_atv_nearest } };
+
+ #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 (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' has not been declared" "" { target c++ } .-1 } */
+ #pragma omp target uses_allocators (memspace(1) : foo) /* { dg-error "expected identifier before numeric constant" } */
+ ; /* { dg-error "expected '\\\)' before ':' token" "" { target c } .-1 } */
+ #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 "'omp_no_such_space' has not been declared" "" { target c++ } .-1 } */
+ #pragma omp target uses_allocators (traits(xyz) : bar) /* { dg-error "traits array must be of 'const omp_alloctrait_t \\\[\\\]' type" "" { target c } } */
+ ; /* { dg-error "'xyz' has not been declared" "" { target c++ } .-1 } */
+ #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" } */
+ ;
+ #pragma omp target uses_allocators (traitz(traits_array), memspace(omp_high_bw_mem_space) : bar) /* { dg-error "unknown modifier 'traitz'" } */
+ ;
+ #pragma omp target uses_allocators (omp_null_allocator) /* { dg-error "'omp_null_allocator' cannot be used in 'uses_allocators' clause" } */
+ ;
+ #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 "legacy 'foo\\\(foo_traits\\\)' traits syntax not allowed in 'uses_allocators' clause when using modifiers" } */
+ ;
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-3.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-3.c
new file mode 100644
index 00000000000..de9ab9213aa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-3.c
@@ -0,0 +1,37 @@
+/* { 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\\(\\)\\) private\\(memspace\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(traits: memspace\\(\\), traits\\(\\)\\) private\\(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\\(\\)\\) private\\(traits\\) private\\(memspace\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(traits: memspace\\(\\), traits\\(mytraits\\)\\) private\\(traits\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace: memspace\\(\\), traits\\(mytraits\\)\\) private\\(memspace\\)" "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/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-4.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-4.c
new file mode 100644
index 00000000000..fabf2a8a319
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-4.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/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_2.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90
index 07327969775..bb984033413 100644
--- a/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90
+++ b/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90
@@ -3,8 +3,6 @@
! Minimal test for valid code:
! - predefined allocators do not need any special treatment in uses_allocators
! (as 'requires dynamic_allocators' is the default).
-!
-! - Non-predefined allocators are currently rejected ('sorry)'
subroutine test
use omp_lib
@@ -35,22 +33,22 @@ subroutine non_predef
integer(kind=omp_allocator_handle_kind) :: a1, a2, a3
- !$omp target uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2)) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+ !$omp target uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2))
block; end block
- !$omp target parallel uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2)) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+ !$omp target parallel uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2))
block; end block
!$omp target uses_allocators(traits(trait):a1) &
- !$omp& uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+ !$omp& uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3)
block; end block
!$omp target parallel uses_allocators(traits(trait):a1) &
- !$omp& uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+ !$omp& uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3)
block; end block
- !$omp target uses_allocators ( traits(trait2) , memspace ( omp_low_lat_mem_space ) : a2 , a3) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+ !$omp target uses_allocators ( traits(trait2) , memspace ( omp_low_lat_mem_space ) : a2 , a3)
block; end block
end subroutine
@@ -62,7 +60,7 @@ subroutine trait_present
integer(kind=omp_allocator_handle_kind) :: a1
! Invalid in OpenMP 5.0 / 5.1, but valid since 5.2 the same as omp_default_mem_space + emptry traits array
- !$omp target uses_allocators ( a1 ) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+ !$omp target uses_allocators ( a1 )
block; end block
end
@@ -76,13 +74,13 @@ subroutine odd_names
integer(kind=omp_allocator_handle_kind) :: traits
integer(kind=omp_allocator_handle_kind) :: memspace
- !$omp target uses_allocators ( traits(trait1), memspace(trait1) ) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+ !$omp target uses_allocators ( traits(trait1), memspace(trait1) )
block; end block
- !$omp target uses_allocators ( traits(trait1), memspace(omp_low_lat_mem_space) : traits) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+ !$omp target uses_allocators ( traits(trait1), memspace(omp_low_lat_mem_space) : traits)
block; end block
- !$omp target uses_allocators ( memspace(omp_low_lat_mem_space), traits(trait1) : memspace) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+ !$omp target uses_allocators ( memspace(omp_low_lat_mem_space), traits(trait1) : memspace)
block; end block
end
@@ -94,6 +92,6 @@ subroutine more_checks
integer(kind=omp_allocator_handle_kind) :: a1, a2(4)
integer(kind=1) :: a3
- !$omp target uses_allocators(memspace (omp_low_lat_mem_space) : a1 ) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+ !$omp target uses_allocators(memspace (omp_low_lat_mem_space) : a1 )
block; end block
end
diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_3.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators_3.f90
new file mode 100644
index 00000000000..8acdd4223f9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/uses_allocators_3.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\\(\\)\\) private\\(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\\)\\) private\\(foo\\) private\\(bar\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(\\), traits\\(traits_array\\)\\) private\\(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\\) private\\(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\\) private\\(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/libgomp/testsuite/libgomp.fortran/uses_allocators_4.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators_4.f90
new file mode 100644
index 00000000000..00f1dcb2763
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/uses_allocators_4.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/libgomp/testsuite/libgomp.fortran/uses_allocators_5.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators_5.f90
new file mode 100644
index 00000000000..00f87109d2c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/uses_allocators_5.f90
@@ -0,0 +1,14 @@
+! { dg-do compile }
+
+program main
+ use omp_lib
+ implicit none
+ integer, allocatable :: arr(:)
+ integer (omp_allocator_handle_kind) :: bar
+
+ !$omp target allocate(bar : arr) ! { dg-error "allocator 'bar' requires 'uses_allocators.bar.' clause in target region" }
+ block
+ allocate(arr(100))
+ end block
+
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_6.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators_6.f90
new file mode 100644
index 00000000000..3799e3cec73
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/uses_allocators_6.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" } }
^ permalink raw reply [flat|nested] 2+ messages in thread
* Re: [patch] OpenMP: Add uses_allocators support
2023-11-20 10:42 [patch] OpenMP: Add uses_allocators support Tobias Burnus
@ 2023-12-08 13:15 ` Jakub Jelinek
0 siblings, 0 replies; 2+ messages in thread
From: Jakub Jelinek @ 2023-12-08 13:15 UTC (permalink / raw)
To: Tobias Burnus; +Cc: gcc-patches, fortran
On Mon, Nov 20, 2023 at 11:42:02AM +0100, Tobias Burnus wrote:
> 2023-11-19 Tobias Burnus <tobias@codesourcery.com>
> Chung-Lin Tang <cltang@codesourcery.com>
>
> gcc/ChangeLog:
>
> * builtin-types.def (BT_FN_VOID_PTRMODE):
> (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,
> BUILT_IN_OMP_DESTROY_ALLOCATOR): Add.
> * omp-low.cc (scan_sharing_clauses):
Missing description.
> +static tree
> +c_parser_omp_clause_uses_allocators (c_parser *parser, tree list)
> +{
> + location_t clause_loc = c_parser_peek_token (parser)->location;
> + tree t = NULL_TREE, nl = list;
> + matching_parens parens;
> + if (!parens.require_open (parser))
> + return list;
> +
> + tree memspace_expr = NULL_TREE;
> + tree traits_var = NULL_TREE;
> +
> + struct item_tok
> + {
> + location_t loc;
> + tree id;
> + item_tok (void) : loc (UNKNOWN_LOCATION), id (NULL_TREE) {}
> + };
> + struct item { item_tok name, arg; };
> + auto_vec<item> *modifiers = NULL, *allocators = NULL;
> + auto_vec<item> *cur_list = new auto_vec<item> (4);
This is certainly the first time I've seen pointers to auto_vec,
normally one uses just vec in such cases, auto_vec is used typically
on automatic variables to make sure the destruction is done.
But I think all the first parse it as a token soup without checking
anything and only in the second round actually check it is something
we've never done before in exactly the same situations.
The usual way would be to quickly peek at tokens to see if there is
: ahead and decide based on that.
See e.g. c_parser_omp_clause_allocate clause.
That has_modifiers check could be basically copied over with
the names of modifiers changed, or could be done in a loop, or
could be moved into a helper function which could be used
by c_parser_omp_clause_allocate and this function and perhaps
others, pass it the list of modifiers and have it return whether
there are modifiers or not.
c_parser_omp_clause_linear does this too (though, that has one of
the modifiers without arguments).
I'm afraid if parsing of every clause does the parsing so significantly
differently it will be a maintainance nightmare.
> @@ -23648,7 +23861,8 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
> | (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)\
Space before \ please (and also in the line above.
> + if (strcmp (IDENTIFIER_POINTER (DECL_NAME (t)),
> + "omp_null_allocator") == 0)
> + {
> + error_at (OMP_CLAUSE_LOCATION (c),
> + "%<omp_null_allocator%> cannot be used in "
> + "%<uses_allocators%> clause");
> + break;
> + }
> +
> + 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 != NULL_TREE
> + && (TREE_CODE (t) != CONST_DECL
> + || 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 "
Maybe %<memspace%> ?
> + "constant enum of %<omp_memspace_handle_t%> type");
> + remove = true;
> + break;
> + }
> + t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);
> + if (t != NULL_TREE)
> + {
> + bool type_err = false;
> +
> + if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE
> + || DECL_SIZE (t) == NULL_TREE)
> + 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 (TREE_CODE (t) != ERROR_MARK)
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;
> + tree cst_val = decl_constant_value_1 (t, true);
> + if (cst_val == t)
> + {
> + error_at (OMP_CLAUSE_LOCATION (c), "traits array must be "
> + "of constant values");
must be initialized with constants?
> + struct item_tok
> + {
> + location_t loc;
> + tree id;
> + item_tok (void) : loc (UNKNOWN_LOCATION), id (NULL_TREE) {}
> + };
> + struct item { item_tok name, arg; };
> + auto_vec<item> *modifiers = NULL, *allocators = NULL;
> + auto_vec<item> *cur_list = new auto_vec<item> (4);
Same as what I wrote in C.
> + case OMP_CLAUSE_USES_ALLOCATORS:
> + t = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
> + 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)
Back to what I wrote about the C++ allocators, if processing_template_decl t
can be type or value dependent and I think we shouldn't stay in the way to
it.
Type dependent as in:
template <typename T, T X>
void foo (T x)
{
#pragma omp target uses_allocators (memspace (...), traits (...): x)
;
#pragma omp target uses_allocators (X)
;
}
instantiated with foo <omp_allocator_handle_t, omp_low_lat_mem_alloc> (...)
or so. So the usual drill, if it is type dependent, defer checking the
type for later. If it is value dependent, don't check exact value and
defer for later.
The above also brings the admittedly preexisting issue, do we want to
support the case where the allocator has type oaht from
typedef omp_allocator_handle_t oaht;
or not? I guess the wording in the standard is unclear on that, users
would certainly expect typedefs/using to work as well (so then
it should be TYPE_MAIN_VARIANT), but perhaps reading OpenMP pedantically
expression of type XYZ is just that type and nothing else.
> + {
> + error_at (OMP_CLAUSE_LOCATION (c),
> + "allocator must be of %<omp_allocator_handle_t%> type");
> + remove = true;
> + break;
> + }
> + if (TREE_CODE (t) == CONST_DECL)
This is another question (also for C, at least C23).
While the allocator needs to be an identifier, so the enumerator or some
variable, not some random other expressions (so e.g. calls to
constexpr/consteval functions are not ok), I wonder whether
constexpr omp_allocator_handle_t myh = omp_low_lat_mem_alloc;
... uses_allocator (myh) ...
is ok or not and ditto the above mentioned uses_allocator (X).
I bet in such case one might not get a CONST_DECL here, and instead if
not value dependent should fold_non_dependent_expr and check if it yields
an INTEGER_CST.
> + {
> + /* Currently for pre-defined allocators in libgomp, we do not
> + require additional init/fini inside target regions, so discard
> + such clauses. */
> + remove = true;
> +
> + if (strcmp (IDENTIFIER_POINTER (DECL_NAME (t)),
> + "omp_null_allocator") == 0)
For omp_null_allocator it would then be integer_zerop check.
> + {
> + error_at (OMP_CLAUSE_LOCATION (c),
> + "%<omp_null_allocator%> cannot be used in "
> + "%<uses_allocators%> clause");
> + break;
> + }
> +
> + 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 != NULL_TREE
> + && (TREE_CODE (t) != CONST_DECL
> + || 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");
I find this OpenMP standard wording just weird but perhaps it is just
because I'm not a native English speaker. I'd write with instead of of.
> + if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE
> + || DECL_SIZE (t) == NULL_TREE)
|| !COMPLETE_TYPE_P (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
Again, question is about typedefs.
> + 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);
> + }
What is unclear to me is how this is supposed to behave in target teams
case. Whether it is acceptable that each team calls omp_init_allocator
and each team has a separate value for the privatized var, or if we'd need
to arrange for it to happen just once and be propagated to other teams
somehow (but that is perhaps impossible).
> + 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);
> + tree t = DECL_INITIAL (traits);
> + gcc_assert (TREE_CODE (t) == CONSTRUCTOR);
> + ntraits = build_int_cst (integer_type_node,
> + CONSTRUCTOR_NELTS (t));
We certainly shouldn't derive ntraits from whether it is a CONSTRUCTOR
and how many CONSTRUCTOR_NELTS it has, it could have fewer than the
array size. IMHO we should go with the array_type_nelts in both cases
(and complain perhaps already in the FEs if it is not constant, we don't
want VLAs there I suppose). Perhaps repeat the CONSTRUCTOR check
for DECL_INITIAL as well earlier when you decide if it should be
firstprivatized or not.
> + t = get_initialized_tmp_var (t, &init_seq, NULL);
> + traits_var = build_fold_addr_expr_loc (loc, t);
> + }
> + else
> + {
> + location_t loc = OMP_CLAUSE_LOCATION (c);
> + gcc_assert (TREE_CODE (TREE_TYPE (traits)) == ARRAY_TYPE);
> + tree t = TYPE_DOMAIN (TREE_TYPE (traits));
> + tree min = TYPE_MIN_VALUE (t);
> + tree max = TYPE_MAX_VALUE (t);
> + gcc_assert (TREE_CODE (min) == INTEGER_CST
> + && TREE_CODE (max) == INTEGER_CST);
> + t = fold_build2_loc (loc, MINUS_EXPR, TREE_TYPE (min),
> + max, min);
> + t = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (min),
> + t, build_int_cst (TREE_TYPE (min), 1));
> + ntraits = t;
Use array_type_nelts.
Jakub
^ permalink raw reply [flat|nested] 2+ messages in thread
end of thread, other threads:[~2023-12-08 13:15 UTC | newest]
Thread overview: 2+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2023-11-20 10:42 [patch] OpenMP: Add uses_allocators support Tobias Burnus
2023-12-08 13:15 ` Jakub Jelinek
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).