From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1410) id B5407388217D; Fri, 30 Jun 2023 19:44:31 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org B5407388217D DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1688154271; bh=BXHn1fc1Wo8bry3dnMRChDrWJADyFusRWOb1ACHlwfk=; h=From:To:Subject:Date:From; b=cavZNHtDNSk4gg4mfXvBJCk3q+1An7zlCpiOyaOj+aAJQrVVnEAwsvxrfFMb2fy5S gEWUtbe6cYYfcLFSlmuvo+tV4/pWU5FKKdcQthqbMb7+tCPmagHZhcLEWCGdFomvtD 139QYfHFqzs5B5s682PluOsMpUhuupd8moabutSE= Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Julian Brown To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-13] OpenMP: OpenMP 5.2 semantics for pointers with unmapped target X-Act-Checkin: gcc X-Git-Author: Julian Brown X-Git-Refname: refs/heads/devel/omp/gcc-13 X-Git-Oldrev: f55d42cc902dfba9a0ae5fd5f670c535afb3f24e X-Git-Newrev: 3acc3ebbbc1529b7207f1c813b3125c947bcf55f Message-Id: <20230630194431.B5407388217D@sourceware.org> Date: Fri, 30 Jun 2023 19:44:31 +0000 (GMT) List-Id: https://gcc.gnu.org/g:3acc3ebbbc1529b7207f1c813b3125c947bcf55f commit 3acc3ebbbc1529b7207f1c813b3125c947bcf55f Author: Julian Brown 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 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 + + * 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 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 + + * 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 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 #include +#include template 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 #include +#include 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 #include +#include 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; };