Hi
On 30/07/2025 1:53 pm, Tobias Burnus wrote:
Can you also update libgomp.texi?
First the sectionhttps://gcc.gnu.org/onlinedocs/libgomp/
OpenMP-5_002e0.html, namely:
"allocate clause" — P — "Clause has no effect on target (PR113436)"
is then 'Y' as the bug is fixed.
Secondly, I wonder whether it makes sense to update
https://gcc.gnu.org/onlinedocs/libgomp/Memory-allocation.html
The main usage is covered by the second bullet point, but maybe we
should add a note
that it also works on 'target' - and, possibly that for firstprivatized
data, there
is first a copy to the device and then to the allocated memory.
I have updated the features table and added a paragraph to describe the
behaviour of allocate in target constructs.
The following testcase that uses a VLA fails with an ICE:
internal compiler error: in force_constant_size, at gimplify.cc:809
use omp_lib
implicit none (type, external)
integer :: s = 7
block
integer :: i(s)
!$omp target firstprivate(i) allocate(allocator(omp_low_lat_mem_alloc) : i)
block
integer :: a(1)
a(i(1)) = 4
end block
end block
end
This type of VLA is expanded to Gimple differently from C-style VLAs. In C:
int x[n]; x[0] = 0;
becomes
int[0:D.3804] * x.28;
int x[0:D.3804] [value-expr: *x.28];
D.3808 = <size calculated from n>;
x.28 = __builtin_alloca_with_align (D.3808, 32);
(*x.28)[0] = 0;
while in Fortran:
integer :: x(n)
x(1) = 1
becomes
integer(kind=4)[0:D.4715] * restrict x;
void * restrict D.4718;
_7 = <size calculated from n>;
D.4718 = __builtin_malloc (_7);
x = D.4718;
(*x)[1] = 1;
In C, the original array is still an array type, with the underlying
dynamically-allocated storage connected to it via it's value-expr. In
Fortran, the array variable is a pointer type and the allocated storage
is directly assigned to it.
In the OpenMP lowering code, I have created a separate function
(is_fortran_variable_sized) to detect this. The code path is mostly
identical to that of C VLAs, but for Fortran there is no need to get the
value-expr and dereference it to get at the backing array. new_var also
needs an indirection before copying for firstprivate as it is a pointer
type rather than an array.
I have made no attempt to optimise for local or USM in this version of
the patch.
Kwok
From c99d1edb73b3146ae7544079b0317942afeca422 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_fortran_variable_sized): New.
(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.
* gfortran.dg/gomp/pr113436-1.f90: New.
* gfortran.dg/gomp/pr113436-2.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 | 249 +++++++++++++++---
gcc/testsuite/c-c++-common/gomp/pr113436-1.c | 29 ++
gcc/testsuite/c-c++-common/gomp/pr113436-2.c | 30 +++
gcc/testsuite/gfortran.dg/gomp/pr113436-1.f90 | 33 +++
gcc/testsuite/gfortran.dg/gomp/pr113436-2.f90 | 36 +++
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 ++++
14 files changed, 703 insertions(+), 42 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/gfortran.dg/gomp/pr113436-1.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/pr113436-2.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..6c6ea8cfa89 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -396,6 +396,20 @@ 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 a Fortran variable sized array. */
+
+static inline bool
+is_fortran_variable_sized (const_tree expr)
+{
+ tree expr_ty = TREE_TYPE (expr);
+ if (lang_GNU_Fortran () && POINTER_TYPE_P (expr_ty))
+ {
+ tree size_unit = TYPE_SIZE_UNIT (TREE_TYPE (expr_ty));
+ return size_unit ? !TREE_CONSTANT (size_unit) : false;
+ }
+ return false;
+}
+
/* Return true if EXPR is variable sized. */
static inline bool
@@ -12813,10 +12827,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 +13017,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 (is_fortran_variable_sized (var))
+ size = TREE_TYPE (size);
+ size = TYPE_SIZE_UNIT (size);
+ if (is_variable_sized (var) || is_fortran_variable_sized (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 +13048,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 +13059,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 +13093,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 (is_fortran_variable_sized (var))
+ size = TREE_TYPE (size);
+ size = TYPE_SIZE_UNIT (size);
+ if (is_variable_sized (var) || is_fortran_variable_sized (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 +13124,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 +14041,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 +14049,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) || is_fortran_variable_sized (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 +14089,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 +14098,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 +14122,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) || is_fortran_variable_sized (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 +14183,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 +14470,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) || is_fortran_variable_sized (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 (!is_fortran_variable_sized (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 +14528,42 @@ 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) || is_fortran_variable_sized (var))
+ {
+ bool fortran_vla_p = is_fortran_variable_sized (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 (!fortran_vla_p)
+ {
+ 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 (fortran_vla_p)
+ 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 +14588,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..0335d025ff7
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/pr113436-1.c
@@ -0,0 +1,29 @@
+/* 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-times "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc
\\\(128, \[0-9\]\+, 5\\\);" 3 "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 "\\\*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..ee762dde6c0
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/pr113436-2.c
@@ -0,0 +1,30 @@
+/* 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-times "D\\\.\[0-9\]+ = __builtin_GOMP_alloc
\\\(64, \[0-9\]+, 4\\\)" 3 "omplower" } } */
+/* { 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/gfortran.dg/gomp/pr113436-1.f90
b/gcc/testsuite/gfortran.dg/gomp/pr113436-1.f90
new file mode 100644
index 00000000000..dc421951dce
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr113436-1.f90
@@ -0,0 +1,33 @@
+! PR middle-end/113436
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-omplower" }
+
+program g
+ use omp_lib
+ implicit none
+
+ integer :: A, B(10), C_arr(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-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..f28fd139bc5
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr113436-2.f90
@@ -0,0 +1,36 @@
+! PR middle-end/113436
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-omplower" }
+
+program g
+ use omp_lib
+ implicit none
+
+ integer :: A, B(10), C_arr(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-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/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
