From patchwork Wed Oct 28 09:43:12 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Jakub Jelinek X-Patchwork-Id: 1389189 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=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=pass (p=none dis=none) header.from=gcc.gnu.org Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.a=rsa-sha256 header.s=default header.b=f8dO7mFr; dkim-atps=neutral Received: from sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4CLkCc4mnKz9sVW for ; Wed, 28 Oct 2020 20:43:24 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 102673857C4D; Wed, 28 Oct 2020 09:43:22 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 102673857C4D DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1603878202; bh=SUVAL8ejexvxzZypP5Hh2DWZEIDJl7vhPWHj9WpgJ+g=; h=Date:To:Subject:List-Id:List-Unsubscribe:List-Archive:List-Post: List-Help:List-Subscribe:From:Reply-To:From; b=f8dO7mFrBaAOWbAvXt2P9J6SNNniA1OseOwJP21/MgPuf1bdoECEihiOTFOYxmHWh 2FHlbraxRN+nizuVKN285sj/Jhv9YArtBsGkS7pA205iFA3QPHCmphKIUjbiZXzhVa BfsScJ1lXHxhqDKXtxyY7yiVzsL1Tujw7BgZCpEk= X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from us-smtp-delivery-124.mimecast.com (us-smtp-delivery-124.mimecast.com [216.205.24.124]) by sourceware.org (Postfix) with ESMTP id 32B6F385801A for ; Wed, 28 Oct 2020 09:43:20 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 32B6F385801A Received: from mimecast-mx01.redhat.com (mimecast-mx01.redhat.com [209.132.183.4]) (Using TLS) by relay.mimecast.com with ESMTP id us-mta-201-ev5tn9jtMOaNTNe2GGcnsA-1; Wed, 28 Oct 2020 05:43:17 -0400 X-MC-Unique: ev5tn9jtMOaNTNe2GGcnsA-1 Received: from smtp.corp.redhat.com (int-mx04.intmail.prod.int.phx2.redhat.com [10.5.11.14]) (using TLSv1.2 with cipher AECDH-AES256-SHA (256/256 bits)) (No client certificate requested) by mimecast-mx01.redhat.com (Postfix) with ESMTPS id E1C3F1074649 for ; Wed, 28 Oct 2020 09:43:16 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-112-37.ams2.redhat.com [10.36.112.37]) by smtp.corp.redhat.com (Postfix) with ESMTPS id 8CF275D9EF for ; Wed, 28 Oct 2020 09:43:16 +0000 (UTC) Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.15.2/8.15.2) with ESMTP id 09S9hD40019278 for ; Wed, 28 Oct 2020 10:43:13 +0100 Received: (from jakub@localhost) by tucnak.zalov.cz (8.15.2/8.15.2/Submit) id 09S9hC09019277 for gcc-patches@gcc.gnu.org; Wed, 28 Oct 2020 10:43:12 +0100 Date: Wed, 28 Oct 2020 10:43:12 +0100 To: gcc-patches@gcc.gnu.org Subject: [committed] openmp: Implicitly discover declare target for variants of declare variant calls Message-ID: <20201028094312.GW7080@tucnak> MIME-Version: 1.0 User-Agent: Mutt/1.11.3 (2019-02-01) X-Scanned-By: MIMEDefang 2.79 on 10.5.11.14 X-Mimecast-Spam-Score: 0 X-Mimecast-Originator: redhat.com Content-Disposition: inline X-Spam-Status: No, score=-6.3 required=5.0 tests=BAYES_00, DKIMWL_WL_HIGH, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, RCVD_IN_DNSWL_NONE, RCVD_IN_MSPIKE_H4, RCVD_IN_MSPIKE_WL, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: , X-Patchwork-Original-From: Jakub Jelinek via Gcc-patches From: Jakub Jelinek Reply-To: Jakub Jelinek Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" Hi! This marks all variants of declare variant also declare target if the base functions are called directly in target regions or declare target functions. Bootstrapped/regtested on x86_64-linux and i686-linux, committed to trunk. 2020-10-28 Jakub Jelinek * omp-offload.c (omp_declare_target_tgt_fn_r): Handle direct calls to declare variant base functions. * testsuite/libgomp.c/target-42.c: New test. Jakub --- gcc/omp-offload.c.jj 2020-10-22 09:26:21.772352421 +0200 +++ gcc/omp-offload.c 2020-10-22 18:54:40.920324767 +0200 @@ -196,7 +196,26 @@ omp_declare_target_var_p (tree decl) static tree omp_discover_declare_target_tgt_fn_r (tree *tp, int *walk_subtrees, void *data) { - if (TREE_CODE (*tp) == FUNCTION_DECL) + if (TREE_CODE (*tp) == CALL_EXPR + && CALL_EXPR_FN (*tp) + && TREE_CODE (CALL_EXPR_FN (*tp)) == ADDR_EXPR + && TREE_CODE (TREE_OPERAND (CALL_EXPR_FN (*tp), 0)) == FUNCTION_DECL + && lookup_attribute ("omp declare variant base", + DECL_ATTRIBUTES (TREE_OPERAND (CALL_EXPR_FN (*tp), + 0)))) + { + tree fn = TREE_OPERAND (CALL_EXPR_FN (*tp), 0); + for (tree attr = DECL_ATTRIBUTES (fn); attr; attr = TREE_CHAIN (attr)) + { + attr = lookup_attribute ("omp declare variant base", attr); + if (attr == NULL_TREE) + break; + tree purpose = TREE_PURPOSE (TREE_VALUE (attr)); + if (TREE_CODE (purpose) == FUNCTION_DECL) + omp_discover_declare_target_tgt_fn_r (&purpose, walk_subtrees, data); + } + } + else if (TREE_CODE (*tp) == FUNCTION_DECL) { tree decl = *tp; tree id = get_identifier ("omp declare target"); @@ -237,7 +256,7 @@ omp_discover_declare_target_tgt_fn_r (tr } if (omp_declare_target_fn_p (decl) || lookup_attribute ("omp declare target host", - DECL_ATTRIBUTES (decl))) + DECL_ATTRIBUTES (decl))) return NULL_TREE; if (!DECL_EXTERNAL (decl) && DECL_SAVED_TREE (decl)) --- libgomp/testsuite/libgomp.c/target-42.c.jj 2020-10-22 18:48:59.679244952 +0200 +++ libgomp/testsuite/libgomp.c/target-42.c 2020-10-22 18:52:51.348903566 +0200 @@ -0,0 +1,42 @@ +#include + +int +on_nvptx (void) +{ + return 1; +} + +int +on_gcn (void) +{ + return 2; +} + +#pragma omp declare variant (on_nvptx) match(construct={target},device={arch(nvptx)}) +#pragma omp declare variant (on_gcn) match(construct={target},device={arch(gcn)}) +int +on (void) +{ + return 0; +} + +int +main () +{ + int v; + #pragma omp target map(from:v) + v = on (); + switch (v) + { + default: + printf ("Host fallback or unknown offloading\n"); + break; + case 1: + printf ("Offloading to NVidia PTX\n"); + break; + case 2: + printf ("Offloading to AMD GCN\n"); + break; + } + return 0; +}