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 <jul...@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. --- gcc/gimplify.cc | 20 ++++++++++++++++++- libgomp/target.c | 7 +++---- .../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 ++++++---- 5 files changed, 40 insertions(+), 14 deletions(-) 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/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; }; -- 2.25.1
