On 13/01/2026 3:06 pm, Tobias Burnus wrote:
Glancing at the code that handles 'allocate' with data sharing clauses,
I see that the following C++ code gives:
#pragma omp parallel firstprivate(aref) allocate(align(128),
allocator(omp_low_lat_mem_alloc) : aref)
→ aref.2 = __builtin_GOMP_alloc (128, 40, 5);
#pragma omp target firstprivate(aref) allocate(align(128),
allocator(omp_low_lat_mem_alloc) : aref)
→ aref.7 = __builtin_GOMP_alloc (128, 8, 5);
The first one (for parallel, old code) uses the size of the array,
the second one (for target, this patch) uses the size of the pointer ...
I think the first one is the expected result.
[This is a C++ program using a reference type]
------------
#include <omp.h>
void f() {
int A[10];
auto &aref = A;
#pragma omp parallel firstprivate(aref) allocate(align(128),
allocator(omp_low_lat_mem_alloc) : aref)
aref[0] = 1;
#pragma omp target firstprivate(aref) allocate(align(128),
allocator(omp_low_lat_mem_alloc) : aref)
aref[0] = 1;
}
-------------
If I look at lower_private_allocate for the C++ program, it is called with:
is_ref=true
The call is in a code block protected by:
else if (omp_privatize_by_reference (var)
And I bet that you should also call that function instead
of using (+ defining) 'is_fortran_variable_sized'.
Example Fortran program:
--------------------
subroutine x
use omp_lib
implicit none
integer :: n = 5
block
integer :: A(n)
!$omp parallel firstprivate(A) allocate(align(128),
allocator(omp_low_lat_mem_alloc) : A)
A(1) = 1;
!$omp end parallel
!$omp target firstprivate(A) allocate(align(128),
allocator(omp_low_lat_mem_alloc) : A)
A(1) = 1;
!$omp end target
end block
end
--------------------
For C++, the hook (→ cxx_omp_privatize_by_reference)
is true for:
return (TYPE_REF_P (TREE_TYPE (decl))
|| is_invisiref_parm (decl));
where the latter is
return ((TREE_CODE (t) == PARM_DECL || TREE_CODE (t) == RESULT_DECL)
&& DECL_BY_REFERENCE (t));
And for Fortran, it is true in a bunch of cases. The hook also
handles, e.g., scalar allocatables or PARM_DECL (reference type or
also with 'optional' attribute).
→ gcc/fortran/trans-openmp.cc:
/* True if OpenMP should privatize what this DECL points to rather
than the DECL itself. */
bool
gfc_omp_privatize_by_reference (const_tree decl)
{
...
I have added an extra is_ref argument to is_variable_sized that acts
similarly to the argument to lower_private_allocate - when true, it
returns true if the object that is referenced is variable sized.
is_variable_sized (var, omp_privatize_by_reference (var)) is now used in
place of is_fortran_variable_sized, and serves to handle Fortran VLAs
and C++ references as well.
Just in case, I have now added checks for the size of memory allocated.
Most common architectures that will be running OpenMP will have
sizeof(int)==4 and sizeof(void*)==8, so I have guarded the checks with
'target int32' or 'target { lp64 || llp64 }'.
I have also added compile-time checks for Fortran statically-allocated
arrays and C++ array references.
* * *
I have made no attempt to optimise for local or USM in this version of
the patch.
It is not completely clear what you mean by 'local'.
I meant when the offloaded code falls back to execution on the host.
Kwok
From 141fe652fa151e6aac20d37fc1db835e23d76868 Mon Sep 17 00:00:00 2001
From: Kwok Cheung Yeung <[email protected]>
Date: Wed, 11 Jun 2025 12:46:44 +0100
Subject: [PATCH] openmp: Allocate memory for private/firstprivate clauses as
directed by allocate clauses in target constructs [PR113436]
This patch generates calls to GOMP_alloc to allocate memory for firstprivate
and private clauses on target constructs with an allocator and alignment
as specified by the allocate clause.
The decl values of the clause need to be adjusted to refer to the allocated
memory, and the initial values of variables need to be copied into the
allocated space for firstprivate variables.
For variable-length arrays, the size of the array is stored in a separate
variable, so the allocation and initialization need to be delayed until the
size is made available on the target.
gcc/
PR middle-end/113436
* omp-low.cc (is_variable_sized): Add extra is_ref argument. Check
referenced type if true.
(lower_omp_target): Call lower_private_allocate to generate code to
allocate memory for firstprivate/private clauses with allocators, and
insert code after dependent variables have been initialized.
Construct calls to free allocate memory and insert after target block.
Adjust decl values for clause variables. Copy value of firstprivate
variables to allocated memory.
gcc/testsuite/
PR middle-end/113436
* c-c++-common/gomp/pr113436-1.c: New.
* c-c++-common/gomp/pr113436-2.c: New.
* g++.dg/gomp/pr113436.C: New.
* gfortran.dg/gomp/pr113436-1.f90: New.
* gfortran.dg/gomp/pr113436-2.f90: New.
* gfortran.dg/gomp/pr113436-3.f90: New.
* gfortran.dg/gomp/pr113436-4.f90: New.
libgomp/
PR middle-end/113436
* libgomp.texi (OpenMP 5.0): Mark allocate clause as implemented.
(Memory allocation): Add documentation for use in target construct.
* testsuite/libgomp.c++/firstprivate-1.C: Enable alignment check.
* testsuite/libgomp.c++/pr113436-1.C: New.
* testsuite/libgomp.c++/pr113436-2.C: New.
* testsuite/libgomp.c++/private-1.C: Enable alignment check.
* testsuite/libgomp.c-c++-common/pr113436-1.c: New.
* testsuite/libgomp.c-c++-common/pr113436-2.c: New.
* testsuite/libgomp.fortran/pr113436-1.f90: New.
* testsuite/libgomp.fortran/pr113436-2.f90: New.
---
gcc/omp-low.cc | 242 +++++++++++++++---
gcc/testsuite/c-c++-common/gomp/pr113436-1.c | 31 +++
gcc/testsuite/c-c++-common/gomp/pr113436-2.c | 32 +++
gcc/testsuite/g++.dg/gomp/pr113436.C | 22 ++
gcc/testsuite/gfortran.dg/gomp/pr113436-1.f90 | 35 +++
gcc/testsuite/gfortran.dg/gomp/pr113436-2.f90 | 38 +++
gcc/testsuite/gfortran.dg/gomp/pr113436-3.f90 | 26 ++
gcc/testsuite/gfortran.dg/gomp/pr113436-4.f90 | 27 ++
libgomp/libgomp.texi | 8 +-
.../testsuite/libgomp.c++/firstprivate-1.C | 6 +-
libgomp/testsuite/libgomp.c++/pr113436-1.C | 27 ++
libgomp/testsuite/libgomp.c++/pr113436-2.C | 25 ++
libgomp/testsuite/libgomp.c++/private-1.C | 3 +-
.../libgomp.c-c++-common/pr113436-1.c | 94 +++++++
.../libgomp.c-c++-common/pr113436-2.c | 80 ++++++
.../testsuite/libgomp.fortran/pr113436-1.f90 | 67 +++++
.../testsuite/libgomp.fortran/pr113436-2.f90 | 58 +++++
17 files changed, 777 insertions(+), 44 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/gomp/pr113436-1.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/pr113436-2.c
create mode 100644 gcc/testsuite/g++.dg/gomp/pr113436.C
create mode 100644 gcc/testsuite/gfortran.dg/gomp/pr113436-1.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/pr113436-2.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/pr113436-3.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/pr113436-4.f90
create mode 100644 libgomp/testsuite/libgomp.c++/pr113436-1.C
create mode 100644 libgomp/testsuite/libgomp.c++/pr113436-2.C
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/pr113436-1.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/pr113436-2.c
create mode 100644 libgomp/testsuite/libgomp.fortran/pr113436-1.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/pr113436-2.f90
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 1444ff4ca43..741aea61128 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -396,11 +396,15 @@ is_taskreg_ctx (omp_context *ctx)
return is_parallel_ctx (ctx) || is_task_ctx (ctx) || is_host_teams_ctx (ctx);
}
-/* Return true if EXPR is variable sized. */
+/* Return true if EXPR is variable sized. If IS_REF is true, then
+ EXPR is assumed to be a reference and the object it refers
+ to is checked instead. */
static inline bool
-is_variable_sized (const_tree expr)
+is_variable_sized (const_tree expr, bool is_ref = false)
{
+ if (is_ref)
+ expr = TREE_TYPE (expr);
return !TREE_CONSTANT (TYPE_SIZE_UNIT (TREE_TYPE (expr)));
}
@@ -12813,10 +12817,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
ilist = NULL;
olist = NULL;
+
+ gimple_seq alloc_dlist = NULL;
+ hash_map<tree, tree> alloc_map;
+ hash_map<tree, gimple_seq> alloc_seq_map;
+
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
{
- tree var, x;
+ tree var, x, new_var, allocator, allocate_ptr, size;
+ gimple_seq alloc_seq;
default:
break;
@@ -12997,10 +13007,27 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
}
map_cnt++;
var = OMP_CLAUSE_DECL (c);
+ new_var = lookup_decl (var, ctx);
+ allocator = NULL_TREE;
+ allocate_ptr = NULL_TREE;
+ size = TREE_TYPE (var);
+ if (omp_privatize_by_reference (var))
+ size = TREE_TYPE (size);
+ size = TYPE_SIZE_UNIT (size);
+ if (is_variable_sized (var, omp_privatize_by_reference (var)))
+ size = lookup_decl (size, ctx);
+ alloc_seq = NULL;
+ if (lower_private_allocate (var, new_var, allocator, allocate_ptr,
+ &alloc_seq, ctx,
+ omp_privatize_by_reference (var),
+ size))
+ {
+ alloc_map.put (new_var, allocate_ptr);
+ alloc_seq_map.put (new_var, alloc_seq);
+ }
if (!omp_privatize_by_reference (var)
&& !is_gimple_reg_type (TREE_TYPE (var)))
{
- tree new_var = lookup_decl (var, ctx);
if (is_variable_sized (var))
{
tree pvar = DECL_VALUE_EXPR (var);
@@ -13011,6 +13038,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
x = build_fold_indirect_ref (new_pvar);
TREE_THIS_NOTRAP (x) = 1;
}
+ else if (allocate_ptr)
+ x = build_fold_indirect_ref (allocate_ptr);
else
x = build_receiver_ref (var, true, ctx);
SET_DECL_VALUE_EXPR (new_var, x);
@@ -13020,6 +13049,26 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR
&& lang_hooks.decls.omp_array_data (var, true))
map_cnt += 2;
+
+ do_dtor:
+ if (allocator)
+ {
+ if (!is_gimple_val (allocator))
+ {
+ tree avar = create_tmp_var (TREE_TYPE (allocator));
+ gimplify_assign (avar, allocator, &alloc_dlist);
+ allocator = avar;
+ }
+ if (!is_gimple_val (allocate_ptr))
+ {
+ tree apvar = create_tmp_var (TREE_TYPE (allocate_ptr));
+ gimplify_assign (apvar, allocate_ptr, &alloc_dlist);
+ allocate_ptr = apvar;
+ }
+ tree f = builtin_decl_explicit (BUILT_IN_GOMP_FREE);
+ gimple *g = gimple_build_call (f, 2, allocate_ptr, allocator);
+ gimple_seq_add_stmt (&alloc_dlist, g);
+ }
break;
case OMP_CLAUSE_PRIVATE:
@@ -13034,7 +13083,25 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
break;
}
var = OMP_CLAUSE_DECL (c);
- if (is_variable_sized (var))
+ new_var = lookup_decl (var, ctx);
+ allocator = NULL_TREE;
+ allocate_ptr = NULL_TREE;
+ alloc_seq = NULL;
+ size = TREE_TYPE (var);
+ if (omp_privatize_by_reference (var))
+ size = TREE_TYPE (size);
+ size = TYPE_SIZE_UNIT (size);
+ if (is_variable_sized (var, omp_privatize_by_reference (var)))
+ size = lookup_decl (size, ctx);
+ lower_private_allocate (var, new_var, allocator, allocate_ptr,
+ &alloc_seq, ctx,
+ omp_privatize_by_reference (var), size);
+ if (allocate_ptr)
+ {
+ alloc_map.put (new_var, allocate_ptr);
+ alloc_seq_map.put (new_var, alloc_seq);
+ }
+ if (!allocate_ptr && is_variable_sized (var))
{
tree new_var = lookup_decl (var, ctx);
tree pvar = DECL_VALUE_EXPR (var);
@@ -13047,7 +13114,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
SET_DECL_VALUE_EXPR (new_var, x);
DECL_HAS_VALUE_EXPR_P (new_var) = 1;
}
- break;
+ goto do_dtor;
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
@@ -13964,7 +14031,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
{
- tree var, x;
+ tree var, x, new_var, *allocate_ptr;
default:
break;
case OMP_CLAUSE_FIRSTPRIVATE:
@@ -13972,10 +14039,32 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
if (is_gimple_omp_oacc (ctx->stmt))
break;
var = OMP_CLAUSE_DECL (c);
+ new_var = lookup_decl (var, ctx);
+ allocate_ptr = alloc_map.get (new_var);
+ if (allocate_ptr)
+ {
+ if (is_variable_sized (var, omp_privatize_by_reference (var)))
+ /* Handle this in the next pass when the size is
+ available. */
+ break;
+
+ gimple_seq *allocate_seq = alloc_seq_map.get (new_var);
+ gcc_assert (allocate_seq);
+ gimple_seq_add_seq (&new_body, *allocate_seq);
+
+ if (omp_privatize_by_reference (var))
+ {
+ x = fold_convert (TREE_TYPE (new_var), *allocate_ptr);
+ gimplify_assign (new_var, x, &new_body);
+ }
+
+ new_var = omp_privatize_by_reference (var)
+ ? build_fold_indirect_ref (new_var)
+ : build_simple_mem_ref (*allocate_ptr);
+ }
if (omp_privatize_by_reference (var)
|| is_gimple_reg_type (TREE_TYPE (var)))
{
- tree new_var = lookup_decl (var, ctx);
tree type;
type = TREE_TYPE (var);
if (omp_privatize_by_reference (var))
@@ -13990,7 +14079,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
x = fold_convert (type, x);
gimplify_expr (&x, &new_body, NULL, is_gimple_val,
fb_rvalue);
- if (omp_privatize_by_reference (var))
+ if (omp_privatize_by_reference (var) && !allocate_ptr)
{
tree v = create_tmp_var_raw (type, get_name (var));
gimple_add_tmp_var (v);
@@ -13999,17 +14088,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
gimple_build_assign (v, x));
x = build_fold_addr_expr (v);
}
- gimple_seq_add_stmt (&new_body,
- gimple_build_assign (new_var, x));
+ gimplify_assign (new_var, x, &new_body);
}
else
{
- bool by_ref = !omp_privatize_by_reference (var);
+ bool by_ref = allocate_ptr
+ || !omp_privatize_by_reference (var);
x = build_receiver_ref (var, by_ref, ctx);
gimplify_expr (&x, &new_body, NULL, is_gimple_val,
fb_rvalue);
- gimple_seq_add_stmt (&new_body,
- gimple_build_assign (new_var, x));
+ gimplify_assign (new_var, x, &new_body);
}
}
else if (is_variable_sized (var))
@@ -14024,29 +14112,59 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
gimple_seq_add_stmt (&new_body,
gimple_build_assign (new_var, x));
}
+ else if (allocate_ptr)
+ {
+ x = build_receiver_ref (var, true, ctx);
+ new_var = unshare_expr (new_var);
+ x = lang_hooks.decls.omp_clause_copy_ctor (c, new_var, x);
+ gimplify_and_add (x, &new_body);
+ }
break;
case OMP_CLAUSE_PRIVATE:
if (is_gimple_omp_oacc (ctx->stmt))
break;
var = OMP_CLAUSE_DECL (c);
+ new_var = lookup_decl (var, ctx);
+ allocate_ptr = alloc_map.get (new_var);
+ if (allocate_ptr)
+ {
+ if (is_variable_sized (var, omp_privatize_by_reference (var)))
+ /* Handle this in the next pass when the size is
+ available. */
+ break;
+
+ gimple_seq *allocate_seq = alloc_seq_map.get (new_var);
+ gcc_assert (allocate_seq);
+ gimple_seq_add_seq (&new_body, *allocate_seq);
+
+ new_var = omp_privatize_by_reference (var)
+ ? new_var
+ : build_simple_mem_ref (*allocate_ptr);
+ }
if (omp_privatize_by_reference (var))
{
location_t clause_loc = OMP_CLAUSE_LOCATION (c);
- tree new_var = lookup_decl (var, ctx);
- x = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_var)));
- if (TREE_CONSTANT (x))
+ if (!allocate_ptr)
{
- x = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_var)),
- get_name (var));
- gimple_add_tmp_var (x);
- TREE_ADDRESSABLE (x) = 1;
- x = build_fold_addr_expr_loc (clause_loc, x);
+ x = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_var)));
+ if (TREE_CONSTANT (x))
+ {
+ x = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_var)),
+ get_name (var));
+ gimple_add_tmp_var (x);
+ TREE_ADDRESSABLE (x) = 1;
+ x = build_fold_addr_expr_loc (clause_loc, x);
+ }
+ else
+ break;
+
+ x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
+ gimplify_expr (&x, &new_body, NULL, is_gimple_val,
+ fb_rvalue);
}
else
- break;
+ x = *allocate_ptr;
- x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
- gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
gimple_seq_add_stmt (&new_body,
gimple_build_assign (new_var, x));
}
@@ -14055,7 +14173,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_HAS_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
- tree new_var;
gimple_seq assign_body;
bool is_array_data;
bool do_optional_check;
@@ -14343,20 +14460,37 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
break;
case OMP_CLAUSE_PRIVATE:
var = OMP_CLAUSE_DECL (c);
- if (is_variable_sized (var))
+ if (is_variable_sized (var, omp_privatize_by_reference (var)))
{
- location_t clause_loc = OMP_CLAUSE_LOCATION (c);
tree new_var = lookup_decl (var, ctx);
- tree pvar = DECL_VALUE_EXPR (var);
- gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
- pvar = TREE_OPERAND (pvar, 0);
+ tree *allocate_ptr = alloc_map.get (new_var);
+ if (allocate_ptr)
+ {
+ gimple_seq *allocate_seq = alloc_seq_map.get (new_var);
+ gcc_assert (allocate_seq);
+ gimple_seq_add_seq (&new_body, *allocate_seq);
+ }
+ location_t clause_loc = OMP_CLAUSE_LOCATION (c);
+ tree pvar = var;
+ if (!omp_privatize_by_reference (var))
+ {
+ pvar = DECL_VALUE_EXPR (var);
+ gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+ pvar = TREE_OPERAND (pvar, 0);
+ }
gcc_assert (DECL_P (pvar));
tree new_pvar = lookup_decl (pvar, ctx);
- tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
- tree al = size_int (DECL_ALIGN (var));
- tree x = TYPE_SIZE_UNIT (TREE_TYPE (new_var));
- x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
- x = fold_convert_loc (clause_loc, TREE_TYPE (new_pvar), x);
+ tree x;
+ if (!allocate_ptr)
+ {
+ tree atmp = builtin_decl_explicit
(BUILT_IN_ALLOCA_WITH_ALIGN);
+ tree al = size_int (DECL_ALIGN (var));
+ x = TYPE_SIZE_UNIT (TREE_TYPE (new_var));
+ x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
+ x = fold_convert_loc (clause_loc, TREE_TYPE (new_pvar), x);
+ }
+ else
+ x = *allocate_ptr;
gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
gimple_seq_add_stmt (&new_body,
gimple_build_assign (new_pvar, x));
@@ -14384,6 +14518,41 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
gimple_build_assign (new_var, x));
}
break;
+ case OMP_CLAUSE_FIRSTPRIVATE:
+ var = OMP_CLAUSE_DECL (c);
+ if (is_variable_sized (var, omp_privatize_by_reference (var)))
+ {
+ tree new_var = lookup_decl (var, ctx);
+ tree *allocate_ptr = alloc_map.get (new_var);
+ if (!allocate_ptr)
+ break;
+ gimple_seq *allocate_seq = alloc_seq_map.get (new_var);
+ gcc_assert (allocate_seq);
+ gimple_seq_add_seq (&new_body, *allocate_seq);
+
+ location_t clause_loc = OMP_CLAUSE_LOCATION (c);
+ tree pvar = var;
+ if (!omp_privatize_by_reference (var))
+ {
+ pvar = DECL_VALUE_EXPR (var);
+ gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+ pvar = TREE_OPERAND (pvar, 0);
+ }
+ gcc_assert (DECL_P (pvar));
+ tree new_pvar = lookup_decl (pvar, ctx);
+ tree x = fold_convert_loc (clause_loc, TREE_TYPE (new_pvar),
+ *allocate_ptr);
+ gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_assign (new_pvar, x));
+
+ x = build_receiver_ref (var, true, ctx);
+ new_var = unshare_expr (new_var);
+ if (omp_privatize_by_reference (var))
+ new_var = build_fold_indirect_ref (new_var);
+ x = lang_hooks.decls.omp_clause_copy_ctor (c, new_var, x);
+ gimplify_and_add (x, &new_body);
+ }
}
gimple_seq fork_seq = NULL;
@@ -14408,6 +14577,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
gimple_seq_add_seq (&new_body, fork_seq);
gimple_seq_add_seq (&new_body, tgt_body);
gimple_seq_add_seq (&new_body, join_seq);
+ gimple_seq_add_seq (&new_body, alloc_dlist);
if (offloaded)
{
diff --git a/gcc/testsuite/c-c++-common/gomp/pr113436-1.c
b/gcc/testsuite/c-c++-common/gomp/pr113436-1.c
new file mode 100644
index 00000000000..985cc212863
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/pr113436-1.c
@@ -0,0 +1,31 @@
+/* PR middle-end/113436 */
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-omplower" } */
+
+#include <omp.h>
+
+void
+f()
+{
+ int A, B[10], *C;
+ A = 5;
+ C = (int *) __builtin_malloc (sizeof (int) * 10);
+ for (int i = 0; i < 10; i++)
+ B[i] = C[i] = i+5;
+
+ #pragma omp target private(A) private(B) private(C)
allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C)
+ {
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+}
+
+/* { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(128,
4, 5\\\);" "omplower" { target int32 } } } */
+/* { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(128,
40, 5\\\);" "omplower" { target int32 } } } */
+/* { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(128,
8, 5\\\);" "omplower" { target { lp64 || llp64 } } } } */
+/* { dg-final { scan-tree-dump "\\\*D\\\.\[0-9\]\+ = 99;" "omplower" } } */
+/* { dg-final { scan-tree-dump "\\\(\\\*D\\\.\[0-9\]\+\\\)\\\[i\\\] =
D\\\.\[0-9\]\+;" "omplower" } } */
+/* { dg-final { scan-tree-dump "\\\*D\\\.\[0-9\]\+ = D\\\.\[0-9\]\+;"
"omplower" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\\(D\\\.\[0-9\]\+,
5\\\);" 3 "omplower" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/pr113436-2.c
b/gcc/testsuite/c-c++-common/gomp/pr113436-2.c
new file mode 100644
index 00000000000..1755b6bd209
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/pr113436-2.c
@@ -0,0 +1,32 @@
+/* PR middle-end/113436 */
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-omplower" } */
+
+#include <omp.h>
+
+void
+g()
+{
+ int A, B[10], *C;
+ A = 5;
+ C = (int *) __builtin_malloc (sizeof (int) * 10);
+ for (int i = 0; i < 10; i++)
+ B[i] = C[i] = i+5;
+
+ #pragma omp target firstprivate(A) firstprivate(B) firstprivate(C)
allocate(allocator(omp_high_bw_mem_alloc), align(64): A, B, C)
+ {
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+}
+
+/* { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(64,
4, 4\\\);" "omplower" { target int32 } } } */
+/* { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(64,
40, 4\\\);" "omplower" { target int32 } } } */
+/* { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(64,
8, 4\\\);" "omplower" { target { lp64 || llp64 } } } } */
+/* { dg-final { scan-tree-dump-times "\\\*D\\\.\[0-9\]\+ = D\\\.\[0-9\]\+;" 3
"omplower" } } */
+/* { dg-final { scan-tree-dump "\\\(\\\*D\\\.\[0-9\]\+\\\) =
\\\(\\\*D\\\.\[0-9\]\+\\\);" "omplower" } } */
+/* { dg-final { scan-tree-dump "\\\*D\\\.\[0-9\]\+ = 99;" "omplower" } } */
+/* { dg-final { scan-tree-dump "\\\(\\\*D\\\.\[0-9\]\+\\\)\\\[i\\\] =
D\\\.\[0-9\]\+;" "omplower" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\\(D\\\.\[0-9\]+,
4\\\)" 3 "omplower" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/pr113436.C
b/gcc/testsuite/g++.dg/gomp/pr113436.C
new file mode 100644
index 00000000000..ad1cc2f83c6
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/pr113436.C
@@ -0,0 +1,22 @@
+// PR middle-end/113436
+// { dg-do "compile" }
+// { dg-options "-std=gnu++20 -fopenmp -fdump-tree-omplower" }
+
+#include <omp.h>
+
+void f()
+{
+ int a[10];
+ auto &aRef = a;
+
+ #pragma omp target firstprivate(aRef) \
+ allocate(align(128), allocator(omp_low_lat_mem_alloc):
aRef)
+ aRef[0] = 1;
+}
+
+// { dg-final { scan-tree-dump "aRef\\\.\[0-9\]\+ = __builtin_GOMP_alloc
\\\(128, 40, 5\\\);" "omplower" { target int32 } } }
+// { dg-final { scan-tree-dump "aRef = aRef\\\.\[0-9\]\+;" "omplower" } }
+// { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = \\\(\\\*D\\\.\[0-9\]+\\\);"
"omplower" } }
+// { dg-final { scan-tree-dump "\\\(\\\*aRef\\\) = D\\\.\[0-9\]\+;" "omplower"
} }
+// { dg-final { scan-tree-dump "\\\(\\\*aRef\\\)\\\[0\\\] = 1;" "omplower" } }
+// { dg-final { scan-tree-dump "__builtin_GOMP_free \\\(aRef\\\.\[0-9\]\+,
5\\\);" "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr113436-1.f90
b/gcc/testsuite/gfortran.dg/gomp/pr113436-1.f90
new file mode 100644
index 00000000000..da757ed720e
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr113436-1.f90
@@ -0,0 +1,35 @@
+! PR middle-end/113436
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-omplower" }
+
+program g
+ use omp_lib
+ implicit none
+
+ integer :: A, B(10)
+ integer, allocatable :: C(:)
+ integer :: i
+
+ A = 5;
+ allocate(C(10))
+ do i = 1, 10
+ B(i) = i + 5
+ C(i) = B(i)
+ end do
+
+ !$omp target private(A) private(B) private(C)
allocate(allocator(omp_high_bw_mem_alloc), align(64): A, B, C)
+ A = 99
+ do i = 1, 10
+ B(i) = -i - 23
+ C(i) = i + 23
+ end do
+ !$omp end target
+end program g
+
+! { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(64,
4, 4\\\);" "omplower" { target int32 } } }
+! { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(64,
40, 4\\\);" "omplower" { target int32 } } }
+! { dg-final { scan-tree-dump-times "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc
\\\(64, \[0-9\]\+, 4\\\);" 3 "omplower" } }
+! { dg-final { scan-tree-dump "\\\*D\\\.\[0-9\]\+ = 99;" "omplower" } }
+! { dg-final { scan-tree-dump
"\\\(\\\*D\\\.\[0-9\]\+\\\)\\\[D\\\.\[0-9\]\+\\\] = D\\\.\[0-9\]\+;" "omplower"
} }
+! { dg-final { scan-tree-dump "D\\\.\[0-9\]\+\\\]\\\[D\\\.\[0-9\]\+\\\] =
D\\\.\[0-9\]\+;" "omplower" } }
+! { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\\(D\\\.\[0-9\]+,
4\\\)" 3 "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr113436-2.f90
b/gcc/testsuite/gfortran.dg/gomp/pr113436-2.f90
new file mode 100644
index 00000000000..0eaf8b5d36d
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr113436-2.f90
@@ -0,0 +1,38 @@
+! PR middle-end/113436
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-omplower" }
+
+program g
+ use omp_lib
+ implicit none
+
+ integer :: A, B(10)
+ integer, allocatable :: C(:)
+ integer :: i
+
+ A = 5;
+ allocate(C(10))
+ do i = 1, 10
+ B(i) = i + 5
+ C(i) = B(i)
+ end do
+
+ !$omp target firstprivate(A) firstprivate(B) firstprivate(C)
allocate(allocator(omp_high_bw_mem_alloc), align(64): A, B, C)
+ A = 99
+ do i = 1, 10
+ B(i) = -i - 23
+ C(i) = i + 23
+ end do
+ !$omp end target
+end program g
+
+! { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(64,
4, 4\\\);" "omplower" { target int32 } } }
+! { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(64,
40, 4\\\);" "omplower" { target int32 } } }
+! { dg-final { scan-tree-dump-times "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc
\\\(64, \[0-9\]\+, 4\\\);" 3 "omplower" } }
+! { dg-final { scan-tree-dump "\\\*D\\\.\[0-9\]\+ = D\\\.\[0-9\]\+;"
"omplower" } }
+! { dg-final { scan-tree-dump "\\\(\\\*D\\\.\[0-9\]\+\\\) =
\\\(\\\*D\\\.\[0-9\]\+\\\);" "omplower" } }
+! { dg-final { scan-tree-dump "\\\*D\\\.\[0-9\]\+ = \\\*D\\\.\[0-9\]\+;"
"omplower" } }
+! { dg-final { scan-tree-dump "\\\*D\\\.\[0-9\]\+ = 99;" "omplower" } }
+! { dg-final { scan-tree-dump
"\\\(\\\*D\\\.\[0-9\]\+\\\)\\\[D\\\.\[0-9\]\+\\\] = D\\\.\[0-9\]\+;" "omplower"
} }
+! { dg-final { scan-tree-dump "D\\\.\[0-9\]\+\\\]\\\[D\\\.\[0-9\]\+\\\] =
D\\\.\[0-9\]\+;" "omplower" } }
+! { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\\(D\\\.\[0-9\]+,
4\\\)" 3 "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr113436-3.f90
b/gcc/testsuite/gfortran.dg/gomp/pr113436-3.f90
new file mode 100644
index 00000000000..f04200f0624
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr113436-3.f90
@@ -0,0 +1,26 @@
+! PR middle-end/113436
+! { dg-do compile }
+! { dg-options "-fopenmp -fno-automatic -fdump-tree-omplower" }
+
+program g
+ use omp_lib
+ implicit none
+
+ integer :: A(10)
+ integer :: i
+
+ do i = 1, 10
+ A(i) = i + 5
+ end do
+
+ !$omp target private(A) allocate(allocator(omp_high_bw_mem_alloc),
align(16): A)
+ do i = 1, 10
+ A(i) = -i + 23
+ end do
+ !$omp end target
+end program g
+
+! { dg-excess-errors "Flag '-fno-automatic' overwrites '-frecursive' implied
by '-fopenmp'" }
+! { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(16,
40, 4\\\);" "omplower" { target int32 } } }
+! { dg-final { scan-tree-dump
"\\\(\\\*D\\\.\[0-9\]\+\\\)\\\[D\\\.\[0-9\]\+\\\] = D\\\.\[0-9\]\+;" "omplower"
} }
+! { dg-final { scan-tree-dump "__builtin_GOMP_free \\\(D\\\.\[0-9\]+, 4\\\)"
"omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr113436-4.f90
b/gcc/testsuite/gfortran.dg/gomp/pr113436-4.f90
new file mode 100644
index 00000000000..19889d21fdf
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr113436-4.f90
@@ -0,0 +1,27 @@
+! PR middle-end/113436
+! { dg-do compile }
+! { dg-options "-fopenmp -fno-automatic -fdump-tree-omplower" }
+
+program g
+ use omp_lib
+ implicit none
+
+ integer :: A(10)
+ integer :: i
+
+ do i = 1, 10
+ A(i) = i + 5
+ end do
+
+ !$omp target firstprivate(A) allocate(allocator(omp_high_bw_mem_alloc),
align(32): A)
+ do i = 1, 10
+ A(i) = -i + 23
+ end do
+ !$omp end target
+end program g
+
+! { dg-excess-errors "Flag '-fno-automatic' overwrites '-frecursive' implied
by '-fopenmp'" }
+! { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(32,
40, 4\\\);" "omplower" { target int32 } } }
+! { dg-final { scan-tree-dump "\\\(\\\*D\\\.\[0-9\]\+\\\) =
\\\(\\\*D\\\.\[0-9\]\+\\\);" "omplower" } }
+! { dg-final { scan-tree-dump
"\\\(\\\*D\\\.\[0-9\]\+\\\)\\\[D\\\.\[0-9\]\+\\\] = D\\\.\[0-9\]\+;" "omplower"
} }
+! { dg-final { scan-tree-dump "__builtin_GOMP_free \\\(D\\\.\[0-9\]+, 4\\\)"
"omplower" } }
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index c6fda669a14..33e585cdef9 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -234,8 +234,7 @@ The OpenMP 4.5 specification is fully supported.
@item Memory management routines @tab Y @tab
@item @code{allocate} directive @tab P
@tab C++ unsupported; see also @ref{Memory allocation}
-@item @code{allocate} clause @tab P @tab Clause has no effect on @code{target}
- (@uref{https://gcc.gnu.org/PR113436,PR113436})
+@item @code{allocate} clause @tab Y @tab
@item @code{use_device_addr} clause on @code{target data} @tab Y @tab
@item @code{ancestor} modifier on @code{device} clause @tab Y @tab
@item Implicit declare target directive @tab Y @tab
@@ -6891,6 +6890,11 @@ The description below applies to:
constant expression with value @code{omp_default_mem_alloc} and no
@code{align} modifier has been specified. (In that case, the normal
@code{malloc} allocation is used.)
+@item The @code{allocate} clause can be used in the @code{target} construct
+ to specify the memory used by @code{private} and @code{firstprivate}
+ variables on offload devices. In the case of @code{firstprivate}, the
+ initial data is first allocated using the default memory allocator, then
+ copied to the memory region specified by the allocator.
@item The @code{allocate} directive for variables in static memory; while
the alignment is honored, the normal static memory is used.
@item Using the @code{allocate} directive for automatic/stack variables, except
diff --git a/libgomp/testsuite/libgomp.c++/firstprivate-1.C
b/libgomp/testsuite/libgomp.c++/firstprivate-1.C
index ae5d4fbe1bf..a7393382cf0 100644
--- a/libgomp/testsuite/libgomp.c++/firstprivate-1.C
+++ b/libgomp/testsuite/libgomp.c++/firstprivate-1.C
@@ -90,14 +90,13 @@ S::g (int dev)
allocate(allocator(omp_low_lat_mem_alloc), align(128):
A, B, C) \
device(dev)
{
-#if 0 /* FIXME: The following is disabled because of PR middle-end/113436. */
if (((uintptr_t) &A) % 128 != 0)
abort ();
if (((uintptr_t) &B) % 128 != 0)
abort ();
if (((uintptr_t) &C) % 128 != 0)
abort ();
-#endif
+
if (A != 5)
abort ();
for (int i = 0; i < 10; i++)
@@ -227,14 +226,13 @@ St<T>::gt (int dev)
allocate(allocator(omp_low_lat_mem_alloc), align(128): A,
B, C) \
device(dev)
{
-#if 0 /* FIXME: The following is disabled because of PR middle-end/113436. */
if (((uintptr_t) &A) % 128 != 0)
abort ();
if (((uintptr_t) &B) % 128 != 0)
abort ();
if (((uintptr_t) &C) % 128 != 0)
abort ();
-#endif
+
if (A != 5)
abort ();
for (int i = 0; i < 10; i++)
diff --git a/libgomp/testsuite/libgomp.c++/pr113436-1.C
b/libgomp/testsuite/libgomp.c++/pr113436-1.C
new file mode 100644
index 00000000000..0aae73b52cf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr113436-1.C
@@ -0,0 +1,27 @@
+/* PR middle-end/113436 */
+/* { dg-do run } */
+
+#include <omp.h>
+#include <stdint.h>
+
+void
+test_int_by_ref ()
+{
+ int a = 5;
+ int &b = a;
+
+ #pragma omp target firstprivate(b) \
+ allocate(allocator(omp_high_bw_mem_alloc), align(64): b)
+ {
+ if (((uintptr_t) &b) % 64 != 0)
+ __builtin_abort ();
+ b *= 7;
+ if (b != 35)
+ __builtin_abort ();
+ }
+}
+
+int main ()
+{
+ test_int_by_ref ();
+}
diff --git a/libgomp/testsuite/libgomp.c++/pr113436-2.C
b/libgomp/testsuite/libgomp.c++/pr113436-2.C
new file mode 100644
index 00000000000..30039950989
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr113436-2.C
@@ -0,0 +1,25 @@
+/* PR middle-end/113436 */
+/* { dg-do run } */
+
+#include <omp.h>
+#include <stdint.h>
+
+void
+test_int_by_ref ()
+{
+ int a = 5;
+ int &b = a;
+
+ #pragma omp target private(b) \
+ allocate(allocator(omp_high_bw_mem_alloc), align(64): b)
+ {
+ if (((uintptr_t) &b) % 64 != 0)
+ __builtin_abort ();
+ b = 7;
+ }
+}
+
+int main ()
+{
+ test_int_by_ref ();
+}
diff --git a/libgomp/testsuite/libgomp.c++/private-1.C
b/libgomp/testsuite/libgomp.c++/private-1.C
index 19ee726a222..84bfc8225f1 100644
--- a/libgomp/testsuite/libgomp.c++/private-1.C
+++ b/libgomp/testsuite/libgomp.c++/private-1.C
@@ -75,14 +75,13 @@ S::g (int dev)
allocate(allocator(omp_low_lat_mem_alloc), align(128): A,
B, C) \
device(dev)
{
-#if 0 /* FIXME: The following is disabled because of PR middle-end/113436. */
if (((uintptr_t) &A) % 128 != 0)
abort ();
if (((uintptr_t) &B) % 128 != 0)
abort ();
if (((uintptr_t) &C) % 128 != 0)
abort ();
-#endif
+
A = 99;
for (int i = 0; i < 10; i++)
B[i] = -i-23;
diff --git a/libgomp/testsuite/libgomp.c-c++-common/pr113436-1.c
b/libgomp/testsuite/libgomp.c-c++-common/pr113436-1.c
new file mode 100644
index 00000000000..18a8792b084
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/pr113436-1.c
@@ -0,0 +1,94 @@
+/* PR middle-end/113436 */
+/* { dg-do run } */
+
+#include <omp.h>
+#include <stdint.h>
+
+void
+test_int_by_val ()
+{
+ int x = 64;
+
+ #pragma omp target firstprivate(x) \
+ allocate(allocator(omp_high_bw_mem_alloc), align(16): x)
+ {
+ if (((uintptr_t) &x) % 16 != 0)
+ __builtin_abort ();
+ x *= 2;
+ if (x != 128)
+ __builtin_abort ();
+ }
+}
+
+void
+test_struct_by_val ()
+{
+ struct S {
+ int a[4];
+ float b[4];
+ } s = { { 1, 2, 3, 4 }, { 5.0f, 6.0f, 7.0f, 8.0f } };
+
+ #pragma omp target firstprivate(s) \
+ allocate(allocator(omp_low_lat_mem_alloc), align(32): s)
+ {
+ if (((uintptr_t) &s) % 32 != 0)
+ __builtin_abort ();
+ for (int i = 0; i < 4; i++)
+ {
+ s.a[i] *= 2;
+ s.b[i] *= 2.0f;
+ }
+ for (int i = 0; i < 4; i++)
+ if (s.a[i] != (i + 1) * 2 || s.b[i] != (i + 5) * 2.0f)
+ __builtin_abort ();
+ }
+}
+
+void
+test_ptr ()
+{
+ int x = 42;
+ int *p = &x;
+ uintptr_t p_orig = (uintptr_t) p;
+ uintptr_t p_new;
+
+ #pragma omp target firstprivate(p) \
+ allocate(allocator(omp_default_mem_alloc), align(16): p) \
+ map(from: p_new)
+ {
+ if (((uintptr_t) &p) % 16 != 0)
+ __builtin_abort ();
+ p_new = (uintptr_t) p;
+ }
+
+ if (p_new != p_orig)
+ __builtin_abort ();
+}
+
+void
+test_vla (int n)
+{
+ int x[n];
+ for (int i = 0; i < n; i++)
+ x[i] = i;
+
+ #pragma omp target firstprivate(x) \
+ allocate(allocator(omp_high_bw_mem_alloc), align(128): x)
+ {
+ if (((uintptr_t) &x) % 128 != 0)
+ __builtin_abort ();
+ for (int i = 0; i < n; i++)
+ x[i]++;
+ for (int i = 0; i < n; i++)
+ if (x[i] != i + 1)
+ __builtin_abort ();
+ }
+}
+
+int main ()
+{
+ test_int_by_val ();
+ test_struct_by_val ();
+ test_ptr ();
+ test_vla (16);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/pr113436-2.c
b/libgomp/testsuite/libgomp.c-c++-common/pr113436-2.c
new file mode 100644
index 00000000000..117944a0e8f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/pr113436-2.c
@@ -0,0 +1,80 @@
+/* PR middle-end/113436 */
+/* { dg-do run } */
+
+#include <omp.h>
+#include <stdint.h>
+
+void
+test_int_by_val ()
+{
+ int x;
+
+ #pragma omp target private(x) \
+ allocate(allocator(omp_high_bw_mem_alloc), align(16): x)
+ {
+ if (((uintptr_t) &x) % 16 != 0)
+ __builtin_abort ();
+ x = 2;
+ }
+}
+
+void
+test_struct_by_val ()
+{
+ struct S {
+ int a[4];
+ float b[4];
+ } s = { { 1, 2, 3, 4 }, { 5.0f, 6.0f, 7.0f, 8.0f } };
+
+ #pragma omp target private(s) \
+ allocate(allocator(omp_low_lat_mem_alloc), align(32): s)
+ {
+ if (((uintptr_t) &s) % 32 != 0)
+ __builtin_abort ();
+ for (int i = 0; i < 4; i++)
+ {
+ s.a[i] = i + 1;
+ s.b[i] = 2.0f * i;
+ }
+ }
+}
+
+void
+test_ptr ()
+{
+ int x = 42;
+ int *p = &x;
+
+ #pragma omp target firstprivate(p) \
+ allocate(allocator(omp_default_mem_alloc), align(16): p)
+ {
+ if (((uintptr_t) &p) % 16 != 0)
+ __builtin_abort ();
+ p++;
+ }
+}
+
+void
+test_vla (int n)
+{
+ int x[n];
+ for (int i = 0; i < n; i++)
+ x[i] = i;
+
+ #pragma omp target private(x) \
+ allocate(allocator(omp_high_bw_mem_alloc), align(128): x)
+ {
+ if (((uintptr_t) &x) % 128 != 0)
+ __builtin_abort ();
+ for (int i = 0; i < n; i++)
+ x[i] = i * 2;
+ }
+}
+
+int main ()
+{
+ test_int_by_val ();
+ test_struct_by_val ();
+ test_ptr ();
+ test_vla (32);
+}
diff --git a/libgomp/testsuite/libgomp.fortran/pr113436-1.f90
b/libgomp/testsuite/libgomp.fortran/pr113436-1.f90
new file mode 100644
index 00000000000..0251525f172
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr113436-1.f90
@@ -0,0 +1,67 @@
+! PR middle-end/113436
+! { dg-do run }
+
+program main
+ use omp_lib
+ implicit none
+
+ call test_integer
+ call test_derived_type
+ call test_vla
+contains
+ subroutine test_integer
+ integer :: x = 64
+
+ !$omp target firstprivate(x) &
+ !$omp & allocate(allocator(omp_high_bw_mem_alloc), align(16): x)
+ if (mod (loc (x), 16) /= 0) stop 1
+ x = x * 2
+ if (x /= 128) stop 2
+ !$omp end target
+ end subroutine
+
+ subroutine test_derived_type
+ type :: Ty
+ integer :: a(4)
+ real*4 :: b(4)
+ end type
+ type (Ty) :: t = Ty (a=(/1, 2, 3, 4/), b=(/5.0, 6.0, 7.0, 8.0/))
+ integer :: i
+
+ !$omp target firstprivate(t) &
+ !$omp & allocate(allocator(omp_low_lat_mem_alloc), align(32): t)
+ if (mod (loc (t), 32) /= 0) stop 3
+ do i = 1, 4
+ t%a(i) = t%a(i) * 2
+ t%b(i) = t%b(i) * 2.0
+ end do
+ do i = 1, 4
+ if (t%a(i) /= i * 2) stop 4
+ if (t%b(i) /= (i + 4) * 2.0) stop 5
+ end do
+ !$omp end target
+ end subroutine
+
+ subroutine test_vla
+ integer :: n = 10
+ integer :: i
+ block
+ integer :: a(n)
+
+ do i = 1, n
+ a(i) = i * 3
+ end do
+
+ !$omp target firstprivate(a) &
+ !$omp & allocate(allocator(omp_low_lat_mem_alloc), align(64): a)
+ if (mod (loc (a), 64) /= 0) stop 6
+ do i = 1, n
+ a(i) = a(i) * 2
+ end do
+ do i = 1, n
+ if (a(i) /= i * 6) stop 7
+ end do
+ !$omp end target
+ end block
+ end subroutine
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/pr113436-2.f90
b/libgomp/testsuite/libgomp.fortran/pr113436-2.f90
new file mode 100644
index 00000000000..2ab257b75e6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr113436-2.f90
@@ -0,0 +1,58 @@
+! PR middle-end/113436
+! { dg-do run }
+
+program main
+ use omp_lib
+ implicit none
+
+ call test_integer
+ call test_derived_type
+contains
+ subroutine test_integer
+ integer :: x
+
+ !$omp target private(x) &
+ !$omp & allocate(allocator(omp_high_bw_mem_alloc), align(16): x)
+ if (mod (loc (x), 16) /= 0) stop 1
+ x = 2
+ !$omp end target
+ end subroutine
+
+ subroutine test_derived_type
+ type :: Ty
+ integer :: a(4)
+ real*4 :: b(4)
+ end type
+ type (Ty) :: t
+ integer :: i
+
+ !$omp target private(t) &
+ !$omp & allocate(allocator(omp_low_lat_mem_alloc), align(32): t)
+ if (mod (loc (t), 32) /= 0) stop 2
+ do i = 1, 4
+ t%a(i) = i
+ t%b(i) = i * 2.0
+ end do
+ !$omp end target
+ end subroutine
+
+ subroutine test_vla
+ integer :: n = 10
+ integer :: i
+ block
+ integer :: a(n)
+
+ do i = 1, n
+ a(i) = i * 3
+ end do
+
+ !$omp target firstprivate(a) &
+ !$omp & allocate(allocator(omp_low_lat_mem_alloc), align(64): a)
+ if (mod (loc (a), 64) /= 0) stop 6
+ do i = 1, n
+ a(i) = a(i) * 2
+ end do
+ !$omp end target
+ end block
+ end subroutine
+end program
--
2.43.0
