Currently, GCC accepts an allocate clause (to use a specific memory
allocator and alignment) on the OpenMP target construct, but it has no
effect - memory is always allocated with the defaults.
This patch causes memory for privatized variables (i.e. variables in
private and firstprivate clauses) to be allocated with the specified
allocator and alignment in a similar fashion to how it is done for
parallel constructs, reusing the lower_private_allocate function.
As the allocated memory is addressed via a pointer, references to the
variables in the target code need to be adjusted to refer to it, which
is done by adjusting the DECL_VALUE_EXPR of the version of the variable
in the target region.
For firstprivate variables, the allocated memory needs to be
initialized. For most part this is done using the existing mechanisms
but to a different target. Arrays need an additional copy of their
contents to the allocated region. C++ references do not need to create a
temporary to hold the referred-to object as the allocated memory
fulfills the role already.
VLAs have a non-constant size which is passed in another variable, so
they cannot be allocated until the size variable is available in the
target region. Similarly to how private VLAs are handled, the allocation
and initialisation is delayed until the size variable is set up.
Tested on a x86_64 host with offloading to nvptx. Okay for trunk?
Kwok
From 84adc8bf84974529e5e73d28c7e0abfd7f421364 Mon Sep 17 00:00:00 2001
From: Kwok Cheung Yeung <kcye...@baylibre.com>
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 (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.
libgomp/
PR middle-end/113436
* 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 | 203 +++++++++++++++---
gcc/testsuite/c-c++-common/gomp/pr113436-1.c | 39 ++++
gcc/testsuite/c-c++-common/gomp/pr113436-2.c | 40 ++++
.../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 | 43 ++++
.../testsuite/libgomp.fortran/pr113436-2.f90 | 38 ++++
11 files changed, 563 insertions(+), 35 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 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 e1036adab28..8efe7c5d2ab 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -12752,10 +12752,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;
@@ -12931,10 +12937,24 @@ 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 = TYPE_SIZE_UNIT (TREE_TYPE (var));
+ if (is_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);
@@ -12945,6 +12965,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);
@@ -12954,6 +12976,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:
@@ -12968,7 +13010,22 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
break;
}
var = OMP_CLAUSE_DECL (c);
+ new_var = lookup_decl (var, ctx);
+ allocator = NULL_TREE;
+ allocate_ptr = NULL_TREE;
+ alloc_seq = NULL;
+ size = TYPE_SIZE_UNIT (TREE_TYPE (var));
if (is_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);
@@ -12981,7 +13038,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:
@@ -13871,7 +13928,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:
@@ -13879,10 +13936,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))
+ /* 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))
@@ -13897,7 +13976,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);
@@ -13906,17 +13985,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))
@@ -13931,29 +14009,54 @@ 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 && !is_variable_sized (var))
+ {
+ 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));
}
@@ -13962,7 +14065,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;
@@ -14252,18 +14354,31 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
var = OMP_CLAUSE_DECL (c);
if (is_variable_sized (var))
{
- location_t clause_loc = OMP_CLAUSE_LOCATION (c);
tree new_var = lookup_decl (var, ctx);
+ 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 = 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));
@@ -14291,6 +14406,35 @@ 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))
+ {
+ 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 = 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);
+ x = lang_hooks.decls.omp_clause_copy_ctor (c, new_var, x);
+ gimplify_and_add (x, &new_body);
+ }
}
gimple_seq fork_seq = NULL;
@@ -14315,6 +14459,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..b78a7cbd292
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/pr113436-1.c
@@ -0,0 +1,39 @@
+/* PR middle-end/113436 */
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-omplower" } */
+
+#include <omp.h>
+#include <stdint.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)
+ {
+ if (((uintptr_t) &A) % 128 != 0)
+ __builtin_abort ();
+ if (((uintptr_t) &B) % 128 != 0)
+ __builtin_abort ();
+ if (((uintptr_t) &C) % 128 != 0)
+ __builtin_abort ();
+ 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 "A\\\.1 = \\\(\[a-z \]*unsigned int\\\)
D\\\.\[0-9\]\+;" "omplower" } } */
+/* { dg-final { scan-tree-dump "B\\\.2 = \\\(\[a-z \]*unsigned int\\\)
D\\\.\[0-9\]\+;" "omplower" } } */
+/* { dg-final { scan-tree-dump "C\\\.3 = \\\(\[a-z \]*unsigned int\\\)
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 "\\\*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..9ad95cc695b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/pr113436-2.c
@@ -0,0 +1,40 @@
+/* PR middle-end/113436 */
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-omplower" } */
+
+#include <omp.h>
+#include <stdint.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)
+ {
+ if (((uintptr_t) &A) % 64 != 0)
+ __builtin_abort ();
+ if (((uintptr_t) &B) % 64 != 0)
+ __builtin_abort ();
+ if (((uintptr_t) &C) % 64 != 0)
+ __builtin_abort ();
+ 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 "A\\\.1 = \\\(\[a-z \]*unsigned int\\\)
D\\\.\[0-9\]+" "omplower" } } */
+/* { dg-final { scan-tree-dump "B\\\.2 = \\\(\[a-z \]*unsigned int\\\)
D\\\.\[0-9\]+" "omplower" } } */
+/* { dg-final { scan-tree-dump "C\\\.3 = \\\(\[a-z \]*unsigned int\\\)
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/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..550903d9843
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr113436-1.f90
@@ -0,0 +1,43 @@
+! 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 = 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
+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..150e42a2524
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr113436-2.f90
@@ -0,0 +1,38 @@
+! 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
+end program
--
2.43.0
