From patchwork Fri Jan 28 16:33:29 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Kwok Cheung Yeung X-Patchwork-Id: 50532 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id B8D21394849B for ; Fri, 28 Jan 2022 16:33:52 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 919D9385840C for ; Fri, 28 Jan 2022 16:33:35 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 919D9385840C Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: zNEty/g/KYb0IibYcUy2uQJCaB7T9sp4IpVY/B6gQnMaxt84rI3mN8pPiSJ36lWkgbtxuYU9hH pEq32Gt0P7513CGgcFfVUzHt0c+Su5Uvlj9U9e3lLivOf7Hfjx7krRGvmSnTxG3MMDHud+URNx zS7zeGTBsGzlokVvtA910A2BY4hFyrX7RVGs7tDu7uSpHVriIyBlfphp3UAYDDdrjfvCa4xvax rd8gOnTAncqZjGWA85QJLsPdESj3oa599bSA9ABD3G0cDOQcMvTmz0BrLNa9ls4i2v7LemrLqX PJpFQcT4T7RhOzHfA1Nibwoz X-IronPort-AV: E=Sophos;i="5.88,324,1635235200"; d="scan'208,223";a="71185662" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 28 Jan 2022 08:33:34 -0800 IronPort-SDR: aXV8hOClQMu7wUs79W+0vpGXkkftD16Q1QRv/H74G85GV4aQ1oNfv5OVuCAr45ySj0QGt+yVvR DOJrzpe0RvvICiN58x+td6rUoHxW0IEmRModjD1hJcp49toMTBwKbe61vi9MagoLNFSI6gMFpC wwA1YWjz2LIF81F09qfHEdhtgBQiMit0QFp0s6rToKpfbFWx6Z7558s75plgGdOdFyvt4aVF2M MyWeE64Xr5Hzm8hTlVhhbOC4bP9P94XYXYXEBB7pixCYL9S8tgelaccpV1PVUzO4lzsOys5qz/ kjs= Message-ID: <2c37dd8b-4a59-6b18-b49e-4c914f45f4f7@codesourcery.com> Date: Fri, 28 Jan 2022 16:33:29 +0000 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (Windows NT 10.0; Win64; x64; rv:91.0) Gecko/20100101 Thunderbird/91.5.1 From: Kwok Cheung Yeung Subject: [PATCH] openmp: Add warning when functions containing metadirectives with 'construct={target}' called directly To: Jakub Jelinek References: <8d413974-0068-3a31-6ae5-d36c1be06d06@codesourcery.com> <20210726142902.GW2380545@tucnak> <0000a35d-75a5-3067-b59c-b4f5bde0ea58@codesourcery.com> <20210726195654.GZ2380545@tucnak> <20210726212329.GE2380545@tucnak> In-Reply-To: <20210726212329.GE2380545@tucnak> X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-08.mgc.mentorg.com (139.181.222.8) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-12.0 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Cc: GCC Patches Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" Hello Regarding this issue which we discussed previously - I have created a patch that adds a warning when this situation is detected. When a metadirective in a explicitly marked target function is gimplified, it is checked to see if it contains a 'construct={target}' selector - if it does, then the containing function is marked with 'omp metadirective construct target'. In the omp-low pass, when function calls are processed, the target function is checked to see if it contains the marker. If it does and the call is not made in a target context, a warning is emitted. This will obviously not catch every possible occurence (e.g. if the function containing the metadirective is called from another target function which is then called locally, or if the call is made via a function pointer), but it might still be useful? Okay for mainline (once the metadirective patches are done)? Thanks Kwok On 26/07/2021 10:23 pm, Jakub Jelinek wrote: > On Mon, Jul 26, 2021 at 10:19:35PM +0100, Kwok Cheung Yeung wrote: >>> Yes, that is a target variant, but I'm pretty sure we've decided that >>> the target construct added for declare target is actually not a dynamic >>> property. So basically mostly return to the 5.0 wording with clarifications >>> for Fortran. See >>> https://github.com/OpenMP/spec/issues/2612#issuecomment-849742988 >>> for details. >>> Making the target in construct dynamic would pretty much force all the >>> scoring to be dynamic as well. >> >> In that comment, Deepak says: >> >> So, we decided to keep the target trait static, requiring that the declare >> target directive must be explicit and that the function version must be >> different from the version of the function that may be called outside of a >> target region (with the additional clarification that whether it differs or >> not will be implementation defined). >> >> "the function version must be different from the version of the function >> that may be called outside of a target region": This is what we do not have >> in GCC at the moment - the function versions called within and outside >> target regions are the same on the host. >> >> "whether it differs or not will be implementation defined": So whether a >> function with 'declare target' and a metadirective involving a 'target' >> construct behaves the same or not when called from both inside and outside >> of a target region is implementation defined? >> >> I will leave the treatment of target constructs in the selector as it is >> then, with both calls going to the same function with the metadirective >> resolving to the 'target' variant. I will try to address your other concerns >> later. > > I think you're right, it should differ in the host vs. target version iff > it is in explicit declare target block, my memory is weak, but let's implement > the 5.0 wording for now (and ignore the 5.1 wording later on) and only when > we'll be doing 5.2 change this (and change for both metadirective and > declare variant at that point). > Ok? > > Jakub > From 741b037a8cd6b85d43a6273ab305ce07705dfa23 Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung Date: Fri, 28 Jan 2022 13:56:33 +0000 Subject: [PATCH] openmp: Add warning when functions containing metadirectives with 'construct={target}' called directly void f(void) { #pragma omp metadirective \ when (construct={target}: A) \ default (B) ... } ... { #pragma omp target f(); // Target call f(); // Local call } With the OpenMP 5.0/5.1 specifications, we would expect A to be selected in the metadirective when the target call is made, but B when f is called directly outside of a target context. However, since GCC does not have separate copies of f for local and target calls, and the construct selector is static, it must be resolved one way or the other at compile-time (currently in the favour of selecting A), which may be unexpected behaviour. This patch attempts to detect the above situation, and will emit a warning if found. 2022-01-28 Kwok Cheung Yeung gcc/ * gimplify.cc (gimplify_omp_metadirective): Mark offloadable functions containing metadirectives with 'construct={target}' in the selector. * omp-general.cc (omp_has_target_constructor_p): New. * omp-general.h (omp_has_target_constructor_p): New prototype. * omp-low.cc (lower_omp_1): Emit warning if marked functions called outside of a target context. gcc/testsuite/ * c-c++-common/gomp/metadirective-4.c (main): Add expected warning. * gfortran.dg/gomp/metadirective-4.f90 (test): Likewise. libgomp/ * testsuite/libgomp.c-c++-common/metadirective-2.c (main): Add expected warning. * testsuite/libgomp.fortran/metadirective-2.f90 (test): Likewise. --- gcc/gimplify.cc | 21 +++++++++++++++++++ gcc/omp-general.cc | 21 +++++++++++++++++++ gcc/omp-general.h | 1 + gcc/omp-low.cc | 18 ++++++++++++++++ .../c-c++-common/gomp/metadirective-4.c | 2 +- .../gfortran.dg/gomp/metadirective-4.f90 | 2 +- .../libgomp.c-c++-common/metadirective-2.c | 2 +- .../libgomp.fortran/metadirective-2.f90 | 2 +- 8 files changed, 65 insertions(+), 4 deletions(-) diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 78bae567ae4..c8a01a4ca52 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -14775,6 +14775,27 @@ gimplify_omp_metadirective (tree *expr_p, gimple_seq *pre_p, gimple_seq *, { auto_vec selectors; + /* Mark offloadable functions containing metadirectives that specify + a 'construct' selector with a 'target' constructor. */ + if (offloading_function_p (current_function_decl)) + { + for (tree clause = OMP_METADIRECTIVE_CLAUSES (*expr_p); + clause != NULL_TREE; clause = TREE_CHAIN (clause)) + { + tree selector = TREE_PURPOSE (clause); + + if (omp_has_target_constructor_p (selector)) + { + tree id = get_identifier ("omp metadirective construct target"); + + DECL_ATTRIBUTES (current_function_decl) + = tree_cons (id, NULL_TREE, + DECL_ATTRIBUTES (current_function_decl)); + break; + } + } + } + /* Try to resolve the metadirective. */ vec candidates = omp_resolve_metadirective (*expr_p); diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc index 4edeb58cc73..842e9fd868f 100644 --- a/gcc/omp-general.cc +++ b/gcc/omp-general.cc @@ -2922,6 +2922,27 @@ omp_resolve_metadirective (gimple *gs) return omp_get_dynamic_candidates (variants); } +bool +omp_has_target_constructor_p (tree selector) +{ + if (selector == NULL_TREE) + return false; + + tree selector_set = TREE_PURPOSE (selector); + if (strcmp (IDENTIFIER_POINTER (selector_set), "construct") != 0) + return false; + + enum tree_code constructs[5]; + int nconstructs + = omp_constructor_traits_to_codes (TREE_VALUE (selector), constructs); + + for (int i = 0; i < nconstructs; i++) + if (constructs[i] == OMP_TARGET) + return true; + + return false; +} + /* Encode an oacc launch argument. This matches the GOMP_LAUNCH_PACK macro on gomp-constants.h. We do not check for overflow. */ diff --git a/gcc/omp-general.h b/gcc/omp-general.h index fdde4a3dfb0..ccdea015e15 100644 --- a/gcc/omp-general.h +++ b/gcc/omp-general.h @@ -126,6 +126,7 @@ extern tree omp_get_context_selector (tree, const char *, const char *); extern tree omp_resolve_declare_variant (tree); extern vec omp_resolve_metadirective (tree); extern vec omp_resolve_metadirective (gimple *); +extern bool omp_has_target_constructor_p (tree); extern tree oacc_launch_pack (unsigned code, tree device, unsigned op); extern tree oacc_replace_fn_attrib_attr (tree attribs, tree dims); extern void oacc_replace_fn_attrib (tree fn, tree dims); diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 59804300c28..07613362ef0 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -14300,6 +14300,24 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) tree fndecl; call_stmt = as_a (stmt); fndecl = gimple_call_fndecl (call_stmt); + if (fndecl + && lookup_attribute ("omp metadirective construct target", + DECL_ATTRIBUTES (fndecl))) + { + bool in_target_ctx = false; + + for (omp_context *up = ctx; up; up = up->outer) + if (gimple_code (up->stmt) == GIMPLE_OMP_TARGET) + { + in_target_ctx = true; + break; + } + if (!ctx || !in_target_ctx) + warning_at (gimple_location (stmt), 0, + "direct calls to an offloadable function containing " + "metadirectives with a % " + "selector may produce unexpected results"); + } if (fndecl && fndecl_built_in_p (fndecl, BUILT_IN_NORMAL)) switch (DECL_FUNCTION_CODE (fndecl)) diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-4.c b/gcc/testsuite/c-c++-common/gomp/metadirective-4.c index c4b109295db..25efbe046bf 100644 --- a/gcc/testsuite/c-c++-common/gomp/metadirective-4.c +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-4.c @@ -25,7 +25,7 @@ void f(double a[], double x) { /* TODO: This does not execute a version of f with the default clause active as might be expected. */ - f (a, 2.71828); + f (a, 2.71828); /* { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" } */ return 0; } diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 index b82c9ea96d9..65eb05cd2fb 100644 --- a/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 @@ -13,7 +13,7 @@ program test ! TODO: This does not execute a version of f with the default clause ! active as might be expected. - call f (a, 2.71828) + call f (a, 2.71828) ! { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" } contains subroutine f (a, x) integer :: i diff --git a/libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c b/libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c index cd5c6c5e21a..55a6098e525 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c +++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c @@ -31,7 +31,7 @@ void f(double a[], double x) { /* TODO: This does not execute a version of f with the default clause active as might be expected. */ - f (a, M_E); + f (a, M_E); /* { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" } */ for (i = 0; i < N; i++) if (fabs (a[i] - (M_E * i)) > EPSILON) diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-2.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-2.f90 index 32017a00077..d83474cf2db 100644 --- a/libgomp/testsuite/libgomp.fortran/metadirective-2.f90 +++ b/libgomp/testsuite/libgomp.fortran/metadirective-2.f90 @@ -19,7 +19,7 @@ program test ! TODO: This does not execute a version of f with the default clause ! active as might be expected. - call f (a, E_CONST) + call f (a, E_CONST) ! { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" } do i = 1, N if (abs (a(i) - (E_CONST * i)) .gt. EPSILON) stop 2