From patchwork Tue Nov 29 15:56:21 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Paul-Antoine Arras X-Patchwork-Id: 27310 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:adf:f944:0:0:0:0:0 with SMTP id q4csp422168wrr; Tue, 29 Nov 2022 07:57:02 -0800 (PST) X-Google-Smtp-Source: AA0mqf73OPeZWvSuSGV2NY+AYOjx2KqQLPtAkDtJjEcWfEgysOpWSRkuhdN0Aer1bolVoPZ6DVSP X-Received: by 2002:a17:906:5e50:b0:7bf:22b6:d18 with SMTP id b16-20020a1709065e5000b007bf22b60d18mr11056037eju.148.1669737422395; Tue, 29 Nov 2022 07:57:02 -0800 (PST) ARC-Seal: i=1; a=rsa-sha256; t=1669737422; cv=none; d=google.com; s=arc-20160816; b=i0tszLUcH6a0VqRCj4I5lvzgIA1Tro5zEBgEQGip6fvydqa0YYY01hmFJhRE+15rV6 TnARy69G2MSiRXwZXhVx9ZF9EqqJr50pKWdCZ3auWicCePUHzKRAcM76lW1rp8SCghts G0vCRbIEqytTElQqFcqsT2tRYtQBKQbavuZl4WLe5DcW9oRFMSrImtJHII4qhEyPPX2D pVExZpQnOpdpFd4/P/KLTJV6IJwJh30lSNf2p7tlkVNHH9+zi6U6DHFaFdkHK5rOKbQb 0EJIcmhNzL2RG9tedHspNlh/DHfNE2qIcVDn/Z+DOEdgfoONOSNTAsljG9I0j1y0xNKK 4yjQ== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:subject:from:to :content-language:user-agent:mime-version:date:message-id :ironport-sdr:dmarc-filter:delivered-to; bh=CblS7KfmOT1YJlcG5B2gWav2fOkSs9d9mUuM7VAwyyM=; b=Mep7ndDOfB8BCUx90UQo7gp30QQ8uifDFBMCrp73uxnv6msELRapofNctJnMmqZ91n mVH1J0qFeXPu1d+qG8tqplS6uM7qnSLVa5B0NAqQRgIrQFmvBuISrH09LHqdSc3bVs1S lIfOSFS9AdMdxuCEMgAIvwY/xzj5fs1wSBiyXmg61Dl+QlSbAwbw/w3Xqwq/s43TU7sx GOE2sXAKCTCl+vWdyAUrg/Sm7dY/n1aP70y77OHpJsYjSwVUo7f6GRKFWoi3+bcq6k9W lxuAJQTl/tbvSnJglzdc9F3uhaLzdu+eWNGXfAC3IUZxlCxcvqvV3jOXT3FQAD++ljWh svRw== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from sourceware.org (server2.sourceware.org. [8.43.85.97]) by mx.google.com with ESMTPS id k7-20020a05640212c700b0046a57695aacsi12385285edx.292.2022.11.29.07.57.02 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Tue, 29 Nov 2022 07:57:02 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) client-ip=8.43.85.97; Authentication-Results: mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 33399385B1B4 for ; Tue, 29 Nov 2022 15:56:55 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id 0664E3858D1E for ; Tue, 29 Nov 2022 15:56:28 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 0664E3858D1E Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.96,203,1665475200"; d="scan'208,223";a="91170238" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 29 Nov 2022 07:56:27 -0800 IronPort-SDR: /1cA3ilzlMh3sEzydd/C6YCDFsV7oOoOBp6QylS5XjgML8Y7V6tdMg0CCEzhaAUjfJarhYEyV/ KMB7drr2br4jeEfKWYUPk0b7uwXbfjQDD+Uk62H7hMw42jXPOk/Xu9FBiyei1vbeXsc4R+XXQD GYA3R7x6d0tdNpLJp94eMSirSTRfRt+PWTx6bd459p8RWXPr7XtAhovXJxjyoDBkysDZwz2Pn7 3izSHmDiw1SYvFSPN7fTyjDLdXaR+yQUHCLL5G+4gv3Gvhr/9o3el37ishfFcBd+VkyVqQkN/N D4E= Message-ID: Date: Tue, 29 Nov 2022 16:56:21 +0100 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (Windows NT 10.0; Win64; x64; rv:102.0) Gecko/20100101 Thunderbird/102.5.0 Content-Language: en-GB To: From: Paul-Antoine Arras Subject: [PATCH] amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) To svr-ies-mbx-13.mgc.mentorg.com (139.181.222.13) X-Spam-Status: No, score=-11.3 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) 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: , Errors-To: gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: =?utf-8?q?INBOX?= X-GMAIL-THRID: =?utf-8?q?1750846587716159800?= X-GMAIL-MSGID: =?utf-8?q?1750846587716159800?= Hi all, This patch adds support for 'gfx803' as an alias for 'fiji' in OpenMP context selectors, so as to be consistent with LLVM. It also adds test cases checking all supported AMD ISAs are properly recognised when used in a 'declare variant' construct. Is it OK for mainline? Thanks, diff --git gcc/config/gcn/gcn.cc gcc/config/gcn/gcn.cc index c74fa007a21..39e93aeaeef 100644 --- gcc/config/gcn/gcn.cc +++ gcc/config/gcn/gcn.cc @@ -2985,7 +2985,7 @@ gcn_omp_device_kind_arch_isa (enum omp_device_kind_arch_isa trait, case omp_device_arch: return strcmp (name, "amdgcn") == 0 || strcmp (name, "gcn") == 0; case omp_device_isa: - if (strcmp (name, "fiji") == 0) + if (strcmp (name, "fiji") == 0 || strcmp (name, "gfx803") == 0) return gcn_arch == PROCESSOR_FIJI; if (strcmp (name, "gfx900") == 0) return gcn_arch == PROCESSOR_VEGA10; diff --git gcc/config/gcn/t-omp-device gcc/config/gcn/t-omp-device index 27d36db894b..538624f7ec7 100644 --- gcc/config/gcn/t-omp-device +++ gcc/config/gcn/t-omp-device @@ -1,4 +1,4 @@ omp-device-properties-gcn: $(srcdir)/config/gcn/gcn.cc echo kind: gpu > $@ echo arch: amdgcn gcn >> $@ - echo isa: fiji gfx900 gfx906 gfx908 gfx90a >> $@ + echo isa: fiji gfx803 gfx900 gfx906 gfx908 gfx90a >> $@ diff --git libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c new file mode 100644 index 00000000000..ae2af1cc00c --- /dev/null +++ libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c @@ -0,0 +1,8 @@ +/* { dg-do run { target { offload_target_amdgcn } } } */ +/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#define USE_FIJI_FOR_GFX803 +#include "declare-variant-4.h" + +/* { dg-final { scan-offload-tree-dump "= gfx803 \\(\\);" "optimized" } } */ diff --git libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c new file mode 100644 index 00000000000..e0437a04d65 --- /dev/null +++ libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c @@ -0,0 +1,7 @@ +/* { dg-do run { target { offload_target_amdgcn } } } */ +/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { scan-offload-tree-dump "= gfx803 \\(\\);" "optimized" } } */ diff --git libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c new file mode 100644 index 00000000000..8de03725dec --- /dev/null +++ libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c @@ -0,0 +1,7 @@ +/* { dg-do run { target { offload_target_amdgcn } } } */ +/* { dg-skip-if "gfx900 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx900" } } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { scan-offload-tree-dump "= gfx900 \\(\\);" "optimized" } } */ diff --git libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c new file mode 100644 index 00000000000..be6f193ed3a --- /dev/null +++ libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c @@ -0,0 +1,7 @@ +/* { dg-do run { target { offload_target_amdgcn } } } */ +/* { dg-skip-if "gfx906 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx906" } } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { scan-offload-tree-dump "= gfx906 \\(\\);" "optimized" } } */ diff --git libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c new file mode 100644 index 00000000000..311fad9074d --- /dev/null +++ libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c @@ -0,0 +1,7 @@ +/* { dg-do run { target { offload_target_amdgcn } } } */ +/* { dg-skip-if "gfx908 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx908" } } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { scan-offload-tree-dump "= gfx908 \\(\\);" "optimized" } } */ diff --git libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c new file mode 100644 index 00000000000..96cc14ca0a3 --- /dev/null +++ libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c @@ -0,0 +1,7 @@ +/* { dg-do run { target { offload_target_amdgcn } } } */ +/* { dg-skip-if "gfx90a only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx90a" } } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { scan-offload-tree-dump "= gfx90a \\(\\);" "optimized" } } */ diff --git libgomp/testsuite/libgomp.c/declare-variant-4.h libgomp/testsuite/libgomp.c/declare-variant-4.h new file mode 100644 index 00000000000..2d7c1ef1a5a --- /dev/null +++ libgomp/testsuite/libgomp.c/declare-variant-4.h @@ -0,0 +1,63 @@ +#pragma omp declare target +int +gfx803 (void) +{ + return 0x803; +} + +int +gfx900 (void) +{ + return 0x900; +} + +int +gfx906 (void) +{ + return 0x906; +} + +int +gfx908 (void) +{ + return 0x908; +} + +int +gfx90a (void) +{ + return 0x90a; +} + +#ifdef USE_FIJI_FOR_GFX803 +#pragma omp declare variant(gfx803) match(device = {isa("fiji")}) +#else +#pragma omp declare variant(gfx803) match(device = {isa("gfx803")}) +#endif +#pragma omp declare variant(gfx900) match(device = {isa("gfx900")}) +#pragma omp declare variant(gfx906) match(device = {isa("gfx906")}) +#pragma omp declare variant(gfx908) match(device = {isa("gfx908")}) +#pragma omp declare variant(gfx90a) match(device = {isa("gfx90a")}) +int +f (void) +{ + return 0; +} + +#pragma omp end declare target + +int +main (void) +{ + int v = 0; + +#pragma omp target map(from : v) + v = f (); + + if (v == 0) + __builtin_abort (); + + __builtin_printf ("AMDGCN accelerator: gfx%x\n", v); + + return 0; +}