From patchwork Mon Jan 20 16:32:00 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 1226067 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-517804-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=mentor.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.a=rsa-sha1 header.s=default header.b=Zxv010V1; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 481cdt0gKbz9sSR for ; Tue, 21 Jan 2020 03:32:32 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :reply-to:from:subject:to:cc:message-id:date:mime-version :content-type; q=dns; s=default; b=cX3tmP4boTLw6R3nAY6R1jos4DJoI duwM4XD1Gsrgoh11ymoV2NOn1QAP3rL8f9pLMihsvXswKAw+izBxz5YVeoeRyNn+ 7wMrRF0ASMX3g5xWA26le9ZnNSYpX5c/omVTdEzFz81i0Tq9AcO4ylprJEhfMuWq V2yhWRx7ihNXuU= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :reply-to:from:subject:to:cc:message-id:date:mime-version :content-type; s=default; bh=SvH6YLx0jXYTL09YSR+UDv5b2Wk=; b=Zxv 010V1+8HXbM6GNtqNV4bDVB16Qf//CrzWLa3zRl4U/xqS2XeAsn9WLuRI/Avxrr3 56/v38yW25wuuMn+fdBiyco4+22vM84YcPLQsWDlYBDdkVJL1p3P2ezezzCcCsfj cFGOm/kQXmfm0VCgRMjhCC7+T9Mv6y+3n8T3pGG4= Received: (qmail 6867 invoked by alias); 20 Jan 2020 16:32:25 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 6856 invoked by uid 89); 20 Jan 2020 16:32:25 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-17.6 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.1 spammy=rec, customer, H*UA:Macintosh, H*u:Macintosh X-HELO: esa2.mentor.iphmx.com Received: from esa2.mentor.iphmx.com (HELO esa2.mentor.iphmx.com) (68.232.141.98) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Mon, 20 Jan 2020 16:32:15 +0000 IronPort-SDR: eMicbh/jzSiOarMy+zIeeX3mS78MQb7rADVEXt2yPeTfxeuoni2U5djP85J1bm1tVFRb6Asb3S I/Xh3BxNHXUL3HsyCII5uXJku2pYxaCfZNhGR43qvgNw6GBdxpiKJ14KZEdgRODwv+bqqN6rWW c+oYvlqPAiRL7wLmOF7rMW9Fd4n6IV7dOED8Vdp/3UJyCnwyGL5ATH1Nb8T0Phx8iDqqH/3TbL FPrZf1bvawsttMMUZ6C72jLxas7Q7KlZ18r6TBRM7tVwF5v9iOOw0pRhFqP3qdpN9M9JNz3sV2 x4Y= Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 20 Jan 2020 08:32:13 -0800 IronPort-SDR: 0eLgkbpvpsaBP5plT8+OJOuf8xFQI9MdMJTHZNSdiSzDrw6QwthbY5RCLhMqY1ehWszf0QytxO uXLsRyDyLgnA== Reply-To: From: Chung-Lin Tang Subject: [PATCH, C++, OpenACC/OpenMP] Allow static constexpr fields in mappable types To: Jakub Jelinek , Thomas Schwinge , gcc-patches CC: Catherine Moore Message-ID: <8b6d4f26-d11b-f564-3712-a6fdda4e4c2a@mentor.com> Date: Tue, 21 Jan 2020 00:32:00 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:68.0) Gecko/20100101 Thunderbird/68.4.1 MIME-Version: 1.0 Hi Jakub, Thomas, We had a customer with a C++ program using GPU offloading failing to compile due to the code's extensive use of 'static constexpr' in its many template classes (code was using OpenMP, but OpenACC is no different) While the FE should ensure that no static members should exist for struct/class types that are being mapped to the GPU, 'static constexpr' are completely resolved and folded statically during compile time, so they really shouldn't count. This is a small patch to cp/decl2.c:cp_omp_mappable_type_1() to allow the DECL_DECLARED_CONSTEXPR_P == true case to be mapped, and a g++ testcase. Patch has been tested with no regressions in g++ and libgomp testsuites. Probably not okay for trunk now, okay for stage1? Thanks, Chung-Lin cp/ * decl2.c (cp_omp_mappable_type_1): Allow fields with DECL_DECLARED_CONSTEXPR_P to be mapped. testsuite/ * g++.dg/goacc/static-constexpr-1.C: New test. diff --git a/gcc/cp/decl2.c b/gcc/cp/decl2.c index 042d6fa12df..4f7d9b0ebd4 100644 --- a/gcc/cp/decl2.c +++ b/gcc/cp/decl2.c @@ -1461,7 +1461,10 @@ cp_omp_mappable_type_1 (tree type, bool notes) { tree field; for (field = TYPE_FIELDS (type); field; field = DECL_CHAIN (field)) - if (VAR_P (field)) + if (VAR_P (field) + /* Fields that are 'static constexpr' can be folded away at compile + time, thus does not interfere with mapping. */ + && !DECL_DECLARED_CONSTEXPR_P (field)) { if (notes) inform (DECL_SOURCE_LOCATION (field), diff --git a/gcc/testsuite/g++.dg/goacc/static-constexpr-1.C b/gcc/testsuite/g++.dg/goacc/static-constexpr-1.C new file mode 100644 index 00000000000..2bf69209de4 --- /dev/null +++ b/gcc/testsuite/g++.dg/goacc/static-constexpr-1.C @@ -0,0 +1,16 @@ +// { dg-do compile } + +/* Test that static constexpr members do not interfere with offloading. */ +struct rec +{ + static constexpr int x = 1; + int y, z; +}; + +void foo (rec& r) +{ + #pragma acc parallel copy(r) + { + r.y = r.y = r.x; + } +}