public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-13] OpenMP: OpenMP 5.2 semantics for pointers with unmapped target
@ 2023-06-30 19:44 Julian Brown
0 siblings, 0 replies; only message in thread
From: Julian Brown @ 2023-06-30 19:44 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:3acc3ebbbc1529b7207f1c813b3125c947bcf55f
commit 3acc3ebbbc1529b7207f1c813b3125c947bcf55f
Author: Julian Brown <julian@codesourcery.com>
Date: Mon Jun 26 14:40:14 2023 +0000
OpenMP: OpenMP 5.2 semantics for pointers with unmapped target
This patch fixes two more cases where an unmapped target pointer results
in a null pointer on the target instead of a copy of the host pointer.
The latter behaviour is required by OpenMP 5.2, which is a change from
earlier versions of the standard. This change has already been made in
one place by Tobias's patch here:
https://gcc.gnu.org/pipermail/gcc-patches/2023-June/622018.html
But this patch makes a similar adjustment in other places
(i.e. for GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION and
GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION).
These changes also revealed a problem with DECL_VALUE_EXPR handling in
gimplify.cc, which this patch also fixes.
2023-06-30 Julian Brown <julian@codesourcery.com>
gcc/
* gimplify.cc (gimplify_scan_omp_clauses): Add note about
DECL_VALUE_EXPR handling for struct mapping nodes.
(gimplify_adjust_omp_clauses): Perform DECL_VALUE_EXPR substitution
before DECL_P check.
libgomp/
* target.c (gomp_map_pointer): Modify zero-length array section
pointer handling.
(gomp_attach_pointer): Likewise.
* testsuite/libgomp.c++/target-lambda-1.C: Update for OpenMP 5.2
semantics.
* testsuite/libgomp.c++/target-this-3.C: Likewise.
* testsuite/libgomp.c++/target-this-4.C: Likewise.
Diff:
---
gcc/ChangeLog.omp | 7 +++++++
gcc/gimplify.cc | 20 +++++++++++++++++++-
libgomp/ChangeLog.omp | 10 ++++++++++
libgomp/target.c | 7 +++----
libgomp/testsuite/libgomp.c++/target-lambda-1.C | 5 ++++-
libgomp/testsuite/libgomp.c++/target-this-3.C | 11 +++++++----
libgomp/testsuite/libgomp.c++/target-this-4.C | 11 +++++++----
7 files changed, 57 insertions(+), 14 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 3485a64b64f..b92e9f5b05f 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,10 @@
+2023-06-30 Julian Brown <julian@codesourcery.com>
+
+ * gimplify.cc (gimplify_scan_omp_clauses): Add note about
+ DECL_VALUE_EXPR handling for struct mapping nodes.
+ (gimplify_adjust_omp_clauses): Perform DECL_VALUE_EXPR substitution
+ before DECL_P check.
+
2023-06-20 Andrew Stubbs <ams@codesourcery.com>
Backport from mainline:
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 707a0c046de..0e856b903ec 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -12090,7 +12090,13 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
/* Adding the decl for a struct access: we haven't created
GOMP_MAP_STRUCT nodes yet, so this statement needs to predict
- whether they will be created in gimplify_adjust_omp_clauses. */
+ whether they will be created in gimplify_adjust_omp_clauses.
+ NOTE: Technically we should probably look through DECL_VALUE_EXPR
+ here because something that looks like a DECL_P may actually be a
+ struct access, e.g. variables in a lambda closure
+ (__closure->__foo) or class members (this->foo). Currently in both
+ those cases we map the whole of the containing object (directly in
+ the C++ FE) though, so struct nodes are not created. */
if (c == grp_end
&& addr_tokens[0]->type == STRUCTURE_BASE
&& addr_tokens[0]->u.structure_base_kind == BASE_DECL
@@ -13895,6 +13901,18 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
remove = true;
break;
}
+ /* If we have a DECL_VALUE_EXPR (e.g. this is a class member and/or
+ a variable captured in a lambda closure), look through that now
+ before the DECL_P check below. (A code other than COMPONENT_REF,
+ i.e. INDIRECT_REF, will be a VLA/variable-length array
+ section. A global var may be a variable in a common block. We
+ don't want to do this here for either of those.) */
+ if ((ctx->region_type & ORT_ACC) == 0
+ && DECL_P (decl)
+ && !is_global_var (decl)
+ && DECL_HAS_VALUE_EXPR_P (decl)
+ && TREE_CODE (DECL_VALUE_EXPR (decl)) == COMPONENT_REF)
+ decl = OMP_CLAUSE_DECL (c) = DECL_VALUE_EXPR (decl);
if (TREE_CODE (decl) == TARGET_EXPR)
{
if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL,
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 2ed60a5e14e..26de5d96fce 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,13 @@
+2023-06-30 Julian Brown <julian@codesourcery.com>
+
+ * target.c (gomp_map_pointer): Modify zero-length array section
+ pointer handling.
+ (gomp_attach_pointer): Likewise.
+ * testsuite/libgomp.c++/target-lambda-1.C: Update for OpenMP 5.2
+ semantics.
+ * testsuite/libgomp.c++/target-this-3.C: Likewise.
+ * testsuite/libgomp.c++/target-this-4.C: Likewise.
+
2023-06-22 Tobias Burnus <tobias@codesourcery.com>
Backported from mainline:
diff --git a/libgomp/target.c b/libgomp/target.c
index fbc84c68952..4447675cd16 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -855,7 +855,7 @@ gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq,
if (n == NULL)
{
if (allow_zero_length_array_sections)
- cur_node.tgt_offset = 0;
+ cur_node.tgt_offset = cur_node.host_start;
else if (devicep->is_usm_ptr_func
&& devicep->is_usm_ptr_func ((void*)cur_node.host_start))
cur_node.tgt_offset = cur_node.host_start;
@@ -1023,9 +1023,8 @@ gomp_attach_pointer (struct gomp_device_descr *devicep,
{
if (allow_zero_length_array_sections)
/* When allowing attachment to zero-length array sections, we
- allow attaching to NULL pointers when the target region is not
- mapped. */
- data = 0;
+ copy the host pointer when the target region is not mapped. */
+ data = target;
else
{
gomp_mutex_unlock (&devicep->lock);
diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-1.C b/libgomp/testsuite/libgomp.c++/target-lambda-1.C
index c5acbb8bf30..fa882d09800 100644
--- a/libgomp/testsuite/libgomp.c++/target-lambda-1.C
+++ b/libgomp/testsuite/libgomp.c++/target-lambda-1.C
@@ -2,6 +2,7 @@
#include <cstdlib>
#include <cstring>
+#include <cstdint>
template <typename L>
void
@@ -22,9 +23,11 @@ struct S
auto fn = [=](void) -> bool
{
bool mapped;
+ uintptr_t hostptr = (uintptr_t) ptr;
+ uintptr_t hostiptr = (uintptr_t) iptr;
#pragma omp target map(from:mapped)
{
- mapped = (ptr != NULL && iptr != NULL);
+ mapped = (ptr != (int*) hostptr && iptr != (int*) hostiptr);
if (mapped)
{
for (int i = 0; i < len; i++)
diff --git a/libgomp/testsuite/libgomp.c++/target-this-3.C b/libgomp/testsuite/libgomp.c++/target-this-3.C
index 6049ba8e201..986582430e2 100644
--- a/libgomp/testsuite/libgomp.c++/target-this-3.C
+++ b/libgomp/testsuite/libgomp.c++/target-this-3.C
@@ -2,6 +2,7 @@
#include <stdio.h>
#include <string.h>
+#include <stdint.h>
extern "C" void abort ();
struct S
@@ -15,12 +16,13 @@ struct S
bool set_ptr (int n)
{
bool mapped;
+ uintptr_t hostptr = (uintptr_t) ptr;
#pragma omp target map(from:mapped)
{
- if (ptr != NULL)
+ if (ptr != (int *) hostptr)
for (int i = 0; i < ptr_len; i++)
ptr[i] = n;
- mapped = (ptr != NULL);
+ mapped = (ptr != (int *) hostptr);
}
return mapped;
}
@@ -28,12 +30,13 @@ struct S
bool set_refptr (int n)
{
bool mapped;
+ uintptr_t hostrefptr = (uintptr_t) refptr;
#pragma omp target map(from:mapped)
{
- if (refptr != NULL)
+ if (refptr != (int *) hostrefptr)
for (int i = 0; i < refptr_len; i++)
refptr[i] = n;
- mapped = (refptr != NULL);
+ mapped = (refptr != (int *) hostrefptr);
}
return mapped;
}
diff --git a/libgomp/testsuite/libgomp.c++/target-this-4.C b/libgomp/testsuite/libgomp.c++/target-this-4.C
index f0237c9b6b8..b2a593d03af 100644
--- a/libgomp/testsuite/libgomp.c++/target-this-4.C
+++ b/libgomp/testsuite/libgomp.c++/target-this-4.C
@@ -4,6 +4,7 @@
#include <cstdlib>
#include <cstring>
+#include <cstdint>
struct T
{
@@ -18,12 +19,13 @@ struct T
auto fn = [=](void) -> bool
{
bool mapped;
+ uintptr_t hostptr = (uintptr_t) ptr;
#pragma omp target map(from:mapped)
{
- if (ptr)
+ if (ptr != (int *) hostptr)
for (int i = 0; i < ptr_len; i++)
ptr[i] = n;
- mapped = (ptr != NULL);
+ mapped = (ptr != (int *) hostptr);
}
return mapped;
};
@@ -35,12 +37,13 @@ struct T
auto fn = [=](void) -> bool
{
bool mapped;
+ uintptr_t hostrefptr = (uintptr_t) refptr;
#pragma omp target map(from:mapped)
{
- if (refptr)
+ if (refptr != (int *) hostrefptr)
for (int i = 0; i < refptr_len; i++)
refptr[i] = n;
- mapped = (refptr != NULL);
+ mapped = (refptr != (int *) hostrefptr);
}
return mapped;
};
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2023-06-30 19:44 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2023-06-30 19:44 [gcc/devel/omp/gcc-13] OpenMP: OpenMP 5.2 semantics for pointers with unmapped target Julian Brown
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).