A while ago, support for the 'allocate' clause on 'omp target' was implemented - for predefined allocators, only.
End of last month, support for defining allocators via the uses_allocators(traits(...), memspace(...) : my_alloc_var) clause was added. This commit fixes mixing the two, i.e. using a such-defined allocator in an 'allocate(my_alloc_var : list,of,vars)' clause. Only some minor tweaks were required to make sure that the right variable got used in the region. Committed as Rev. r17-2421-gabe9e1f1698520. (Post-commit) comment are highly welcome, as usual! Thanks, Tobias
commit abe9e1f169852097377e0afca62d88df816d9cac Author: Tobias Burnus <[email protected]> Date: Wed Jul 15 18:43:19 2026 +0200 OpenMP: 'allocate' on 'target' with uses_allocators' def'ed allocator This implements the tweaks required for omp target uses_allocators(traits(...): my_alloc) allocate(my_alloc:...) i.e. using an allocator variable in 'allocate' on target, which is initialized via the 'uses_allocators' directive. ('allocate' on target may either use such an allocator or a predefined allocator. Using a random allocator handle is rejected at compile time.) gcc/ChangeLog: * omp-low.cc (lower_private_allocate): When called on 'target' and the allocator is a variable, obtain the device variable. (lower_omp_target): Move processing of OMP_CLAUSE_USES_ALLOCATORS on the device/receiver side in before the loop over clauses to ensure the VALUE_EXPR points to uses_allocators's inited var. libgomp/ChangeLog: * testsuite/libgomp.c-c++-common/uses_allocators-10.c: New test. gcc/testsuite/ChangeLog: * c-c++-common/gomp/uses_allocators-3.c: New test. --- gcc/omp-low.cc | 27 ++++++--- .../c-c++-common/gomp/uses_allocators-3.c | 65 ++++++++++++++++++++++ .../libgomp.c-c++-common/uses_allocators-10.c | 28 ++++++++++ 3 files changed, 112 insertions(+), 8 deletions(-) diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index ec421657c68..33082fba1a3 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -4884,6 +4884,12 @@ lower_private_allocate (tree var, tree new_var, tree &allocator, return false; } + if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET && DECL_P (allocator)) + { + allocator = lookup_decl (allocator, ctx); + gcc_checking_assert (allocator != NULL_TREE); + } + unsigned HOST_WIDE_INT ialign = 0; if (TREE_CODE (allocator) == TREE_LIST) { @@ -12942,6 +12948,19 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) hash_map<tree, tree> alloc_map; hash_map<tree, gimple_seq> alloc_seq_map; + /* The value expression for USES_ALLOCATORS needs to be setup before the + allocator is used via lower_private_allocate. */ + for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USES_ALLOCATORS) + { + tree allocator = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c); + tree new_allocator = lookup_decl (allocator, ctx); + tree x = build_receiver_ref (allocator, false, ctx); + SET_DECL_VALUE_EXPR (new_allocator, x); + DECL_HAS_VALUE_EXPR_P (new_allocator) = 1; + map_cnt++; + } + for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) switch (OMP_CLAUSE_CODE (c)) { @@ -13287,14 +13306,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) DECL_HAS_VALUE_EXPR_P (new_var) = 1; } break; - case OMP_CLAUSE_USES_ALLOCATORS: - allocator = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c); - tree new_allocator = lookup_decl (allocator, ctx); - x = build_receiver_ref (allocator, false, ctx); - SET_DECL_VALUE_EXPR (new_allocator, x); - DECL_HAS_VALUE_EXPR_P (new_allocator) = 1; - map_cnt++; - break; } if (offloaded) 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..c92e0f56f6c --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-3.c @@ -0,0 +1,65 @@ +// { dg-do compile } + +// Check that the all non-predefined allocators used in +// target's allocator clauses are in 'uses_allocators' + +//#include <omp.h> + +typedef __UINTPTR_TYPE__ omp_uintptr_t; + +typedef enum omp_alloctrait_key_t +{ + omp_atk_alignment = 2, +} omp_alloctrait_key_t; + +typedef enum omp_alloctrait_value_t +{ + omp_atv_default = (__UINTPTR_TYPE__) -1, +} omp_alloctrait_value_t; + +typedef struct omp_alloctrait_t +{ + omp_alloctrait_key_t key; + omp_uintptr_t value; +} omp_alloctrait_t; + +typedef enum omp_allocator_handle_t +{ + omp_null_allocator = 0, + __omp_allocator_handle_t_max__ = __UINTPTR_MAX__ +} omp_allocator_handle_t; + +typedef enum omp_memspace_handle_t +{ + omp_default_mem_space = 0, + omp_low_lat_mem_space = 4, + __omp_memspace_handle_t_max__ = __UINTPTR_MAX__ +} omp_memspace_handle_t; + + +int main () +{ + const omp_alloctrait_t traits[] = {{omp_atk_alignment, 1024}}; + omp_allocator_handle_t alloc = omp_null_allocator; + omp_allocator_handle_t alloc2 = omp_null_allocator; + + int array[2], var2 = 124; + int res = 1; + + #pragma omp target map(from: res) \ + uses_allocators(traits(traits) : alloc2) \ + allocate(alloc : array, var2) private(array) firstprivate(var2) + { +// { dg-error "allocator 'alloc' requires 'uses_allocators\\(alloc\\)' clause in target region" "" { target *-*-* } .-2 } + array[0] = 1; + array[1] = 2; + if (var2 != 124 || ((omp_uintptr_t) &var2) % 1024 != 0) + __builtin_abort (); + else if (array[0] != 1 || array[1] != 2 || ((omp_uintptr_t) &array[0]) % 1024 != 0) + __builtin_abort (); + else + res = 42; + } + if (res != 42) + __builtin_abort (); +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-10.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-10.c new file mode 100644 index 00000000000..95469dcab1e --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-10.c @@ -0,0 +1,28 @@ +// { dg-do run } + +#include <omp.h> + +int main () +{ + const omp_alloctrait_t traits[] = {{omp_atk_alignment, 1024}}; + omp_allocator_handle_t alloc = omp_null_allocator; + + int array[2], var2 = 124; + int res = 1; + + #pragma omp target map(from: res) \ + uses_allocators(memspace(omp_default_mem_space), traits(traits) : alloc) \ + allocate(alloc : array, var2) private(array) firstprivate(var2) + { + array[0] = 1; + array[1] = 2; + if (var2 != 124 || ((omp_uintptr_t) &var2) % 1024 != 0) + __builtin_abort (); + else if (array[0] != 1 || array[1] != 2 || ((omp_uintptr_t) &array[0]) % 1024 != 0) + __builtin_abort (); + else + res = 42; + } + if (res != 42) + __builtin_abort (); +}
