[gcc r17-2421] OpenMP: 'allocate' on 'target' with uses_allocators' def'ed allocator

https://gcc.gnu.org/g:abe9e1f169852097377e0afca62d88df816d9cac

commit r17-2421-gabe9e1f169852097377e0afca62d88df816d9cac
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.

Diff:
---
 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 ec421657c68b..33082fba1a36 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 000000000000..c92e0f56f6cc
--- /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 000000000000..95469dcab1ed
--- /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 ();
+}