From patchwork Wed Nov 30 15:32:06 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: 27857 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:adf:f944:0:0:0:0:0 with SMTP id q4csp995315wrr; Wed, 30 Nov 2022 07:32:49 -0800 (PST) X-Google-Smtp-Source: AA0mqf4jWRw1OeDSULDqnLFgem8iHXANsCnRIJ6R86bWyysNtU4qpRaPxUKIV/YYxIh/i4MJeIoL X-Received: by 2002:a17:906:6c89:b0:7c0:83c2:24a4 with SMTP id s9-20020a1709066c8900b007c083c224a4mr7676827ejr.251.1669822369043; Wed, 30 Nov 2022 07:32:49 -0800 (PST) ARC-Seal: i=1; a=rsa-sha256; t=1669822369; cv=none; d=google.com; s=arc-20160816; b=oAZ+nsd7dwQ4iVPUO8DJd33PagMqq1niKLUm/82p9kOv0ZEkxh2cJPOt/ObPKIp3qg eRwvryUjyxGJAaPMxzhmggTlKrZGGETSNEB5raURWTyMsPChuvHJ6Fx5NEOhEPNLwMlu XDfBua9BXa74b+BOxfEL77pu33JlU/Kiam69zIXfvBtdA2FLuPbQnglznqFevpNt+p9C G1iETl9XjYTpATx65SsATcm6ljIrWzHXCCG15RG1l8ZmbqB7O7mW/V0aU4iJBPvCrviq N9Qim+zkPN/2VF9jxqpZ9FqmwBcOUbYEFOHx7UUqgLF/oDT7Kr2XYAEfkggKKDPRxel8 2/qg== 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=mVbbucmu+Z96ELocAyB7/gEpIbqGDJkufdX0SEKx0GI=; b=pj4aq2kAOX4WvBlPZT+tJv5oF7rSPuksZMHUxqlkGOXIABGl8i1pl5+tHFogNT7mjI po9Y1VnvdsEdWlcbe305kQV79ZeoCuRu27wNHt2mLC/iNrbWWI2CMjnT84zBBKlClRjY 9YZUkOGIJ8h5im0GZUXQzMYv+Ro0TfLs2cxJsjB8ijjC3IcIJcpwJc2u6lWONPnEiyDt gInfD72mm8eyaWpW3ipA8sxuosBZdFKo4PGsqFwnRN8Vm9ZwjiBIU/uZJpQEx/bMZAhe DuNis4UzGHzA2aUmhc520xq8f/2OSXtexWNm4IvAlrj3dEnLc4eRpqHDNpN/tUF2qooQ KuWw== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from sourceware.org (server2.sourceware.org. [2620:52:3:1:0:246e:9693:128c]) by mx.google.com with ESMTPS id sb39-20020a1709076da700b007aeec2bf18esi1687229ejc.529.2022.11.30.07.32.48 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 30 Nov 2022 07:32:49 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) client-ip=2620:52:3:1:0:246e:9693:128c; Authentication-Results: mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c 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 9520038532E2 for ; Wed, 30 Nov 2022 15:32:39 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 1B61E385B193 for ; Wed, 30 Nov 2022 15:32:13 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 1B61E385B193 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,206,1665475200"; d="scan'208,223";a="88589969" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 30 Nov 2022 07:32:11 -0800 IronPort-SDR: OFJ2Oww9meHrHtb3Sx+/zvw0CRO5CWoZhGGxy8FDlc5dhQ+cMaeGV4eShCh9P1+e8FV0wKY0Tv qI8SpEFLGRyP/6BfUdvuqZBjLp9Qs3Cy/EyLsGTojwmg5c1jk/APgZM7usNiZU36nIR3pFQ0Le FayfqLT7fPWPXJBYzJgcg3zMwvDKxnzLoF0RJ8sCEfaYplf5IFBrcgCrwEJZ4JAIi03GbSMs/V JtSGBNW05SFt9Hy/kqL7VgAh5kYYwlYV+zGXuzS1Romqw4mJcZKmkaxkxmTp76Ylg9zXWMrWck 2PQ= Message-ID: Date: Wed, 30 Nov 2022 16:32:06 +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][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-07.mgc.mentorg.com (139.181.222.7) To svr-ies-mbx-13.mgc.mentorg.com (139.181.222.13) X-Spam-Status: No, score=-11.5 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, RCVD_IN_MSPIKE_H2, 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?1750935660211808433?= X-GMAIL-MSGID: =?utf-8?q?1750935660211808433?= Hi all, This patch adds or fixes support for various AMD 'isa' and 'arch' trait 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 'metadirective' construct. This patch is closely related to https://gcc.gnu.org/r13-4403-g1fd508744eccda but cannot be committed to mainline because metadirectives and dynamic context selectors have not landed there yet. Can this be committed to OG12? Thanks, From 88522107dd39ba3ff8465cf688fe4438fa3b77b4 Mon Sep 17 00:00:00 2001 From: Paul-Antoine Arras Date: Wed, 30 Nov 2022 14:52:55 +0100 Subject: [PATCH] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors Add or fix libgomp support for 'amdgcn' as arch, and 'gfx908' and 'gfx90a' as isa traits. Add test case for all supported 'isa' values used as context selectors in a metadirective construct.. libgomp/ChangeLog: * config/gcn/selector.c (GOMP_evaluate_current_device): Recognise 'amdgcn' as arch, and 'gfx908' and 'gfx90a' as isa traits. * testsuite/libgomp.c-c++-common/metadirective-6.c: New test. --- libgomp/config/gcn/selector.c | 15 ++++-- .../libgomp.c-c++-common/metadirective-6.c | 48 +++++++++++++++++++ 2 files changed, 60 insertions(+), 3 deletions(-) create mode 100644 libgomp/testsuite/libgomp.c-c++-common/metadirective-6.c diff --git libgomp/config/gcn/selector.c libgomp/config/gcn/selector.c index 60793fc05d3..c948497c538 100644 --- libgomp/config/gcn/selector.c +++ libgomp/config/gcn/selector.c @@ -36,7 +36,7 @@ GOMP_evaluate_current_device (const char *kind, const char *arch, if (kind && strcmp (kind, "gpu") != 0) return false; - if (arch && strcmp (arch, "gcn") != 0) + if (arch && (strcmp (arch, "gcn") != 0 || strcmp (arch, "amdgcn") != 0)) return false; if (!isa) @@ -48,8 +48,17 @@ GOMP_evaluate_current_device (const char *kind, const char *arch, #endif #ifdef __GCN5__ - if (strcmp (isa, "gfx900") == 0 || strcmp (isa, "gfx906") != 0 - || strcmp (isa, "gfx908") == 0) + if (strcmp (isa, "gfx900") == 0 || strcmp (isa, "gfx906") != 0) + return true; +#endif + +#ifdef __CDNA1__ + if (strcmp (isa, "gfx908") == 0) + return true; +#endif + +#ifdef __CDNA2__ + if (strcmp (isa, "gfx90a") == 0) return true; #endif diff --git libgomp/testsuite/libgomp.c-c++-common/metadirective-6.c libgomp/testsuite/libgomp.c-c++-common/metadirective-6.c new file mode 100644 index 00000000000..6d169001db1 --- /dev/null +++ libgomp/testsuite/libgomp.c-c++-common/metadirective-6.c @@ -0,0 +1,48 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options "-foffload=-fdump-tree-omp_expand_metadirective" } */ + +#define N 100 + +void f (int x[], int y[], int z[]) +{ + int i; + + #pragma omp target map(to: x, y) map(from: z) + #pragma omp metadirective \ + when (device={isa("gfx803")}: teams num_teams(512)) \ + when (device={isa("gfx900")}: teams num_teams(256)) \ + when (device={isa("gfx906")}: teams num_teams(128)) \ + when (device={isa("gfx908")}: teams num_teams(64)) \ + when (device={isa("gfx90a")}: teams num_teams(32)) \ + default (teams num_teams(4)) + for (i = 0; i < N; i++) + z[i] = x[i] * y[i]; +} + +int main (void) +{ + int x[N], y[N], z[N]; + int i; + + for (i = 0; i < N; i++) + { + x[i] = i; + y[i] = -i; + } + + f (x, y, z); + + for (i = 0; i < N; i++) + if (z[i] != x[i] * y[i]) + return 1; + + return 0; +} + +/* The metadirective should be resolved after Gimplification. */ + +/* { dg-final { scan-offload-tree-dump "__builtin_GOMP_teams4 \\(512, 512" "omp_expand_metadirective" { target { any-opts "-foffload=-march=fiji" } } } } */ +/* { dg-final { scan-offload-tree-dump "__builtin_GOMP_teams4 \\(256, 256" "omp_expand_metadirective" { target { any-opts "-foffload=-march=gfx900" } } } } */ +/* { dg-final { scan-offload-tree-dump "__builtin_GOMP_teams4 \\(128, 128" "omp_expand_metadirective" { target { any-opts "-foffload=-march=gfx906" } } } } */ +/* { dg-final { scan-offload-tree-dump "__builtin_GOMP_teams4 \\(64, 64" "omp_expand_metadirective" { target { any-opts "-foffload=-march=gfx908" } } } } */ +/* { dg-final { scan-offload-tree-dump "__builtin_GOMP_teams4 \\(32, 32" "omp_expand_metadirective" { target { any-opts "-foffload=-march=gfx90a" } } } } */