From patchwork Thu Nov 30 15:15:04 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 171952 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a59:bcd1:0:b0:403:3b70:6f57 with SMTP id r17csp460160vqy; Thu, 30 Nov 2023 07:15:38 -0800 (PST) X-Google-Smtp-Source: AGHT+IHVCcDBgZFlk4u0OPzTU9x9Q+EfCMPshJ+a6UzXDwpUqrOLSdM839M+LUvE1W5zEV9eYL8C X-Received: by 2002:a05:6871:4504:b0:1f0:b2a:768a with SMTP id nj4-20020a056871450400b001f00b2a768amr29988568oab.8.1701357338142; Thu, 30 Nov 2023 07:15:38 -0800 (PST) ARC-Seal: i=2; a=rsa-sha256; t=1701357338; cv=pass; d=google.com; s=arc-20160816; b=vrEVb1sAOXDUzFLh21AmnQm/oYNCCbmJw8KMaLGjQwXQcHnwzL1q16n9VXAFMliiJK guIBRprKsteD5eqiiLJrpQPfRRAuJtd+ersP64WxsBI/L8LJK5u1fPq4+0umxw/XlTjJ 2onykLuwXVb/K/l9rSzDT90HdL/kjMAPawh25MaH+tnySYhocm38CTzT2D4RPlovSB9R SSiled7K2y1edXJ4PWAjS0qUKHsnIjS/CoLSGUfSV4hbWF/1zjNYRReVghuR0a2GFKI4 XY3Em2RIAHi7nfWcTBuxP3JdkFZag0c0VL2AWOpguPnPWu8ZWQ8RT0R115XonBdRrkq8 7OBQ== ARC-Message-Signature: i=2; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:mime-version:message-id:date :user-agent:references:in-reply-to:subject:cc:to:from:ironport-sdr :arc-filter:dmarc-filter:delivered-to; bh=EzOxE7lbhg/QOo45aBBt9IkDcXDB0xLROpkGeYH5/eg=; fh=xFwJ5ROgdbV2q0+SyWqec+uG4TktcAEcABK3HYj3ymw=; b=b5e2DSrxuSuXNZHBNiIGhfbLDTzW/m5dRbMSYfghpBStrwnZ68cZ8mRchjIZCWcZDz kbEwHR6kqwZlqBOxXI+ccbQH1mkV6mdFCo4tDo7EXkZQkdWfOeC9BKWqgPybQh3zayUh vV37iA9beXvQeV+m9fOA/7HPZDuxKJ4Ih8N9N3ac1vliMbPnY3NrBabV6vUGQ1S6Jeof bFMMDu/5jTy7r3lS1YLqLltMi3H5jOuTrz5paXAu8/IbtzoQZEcp1bSJtK501QyRztSj vncHEKjFxo3Yt8H2aTfp6Y7zHa/Ucb8RXi2FK+y0cmiE29ZU7D07I10wy5nS8woeoVMx I0YA== ARC-Authentication-Results: i=2; mx.google.com; arc=pass (i=1); 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 (server2.sourceware.org. [2620:52:3:1:0:246e:9693:128c]) by mx.google.com with ESMTPS id v13-20020a056102226d00b00464540a733asi195022vsd.66.2023.11.30.07.15.38 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Thu, 30 Nov 2023 07:15:38 -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; arc=pass (i=1); 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 D7C693858298 for ; Thu, 30 Nov 2023 15:15:37 +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 321FF3858D37 for ; Thu, 30 Nov 2023 15:15:13 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 321FF3858D37 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 321FF3858D37 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=68.232.129.153 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1701357315; cv=none; b=YPCzk4VrtctIUvPm36+4z2cJzYlaXbSirZZzMJY9Wcrvzv+BgzZ+OBre0yVywrqdHM/RL6/O+5mCig9o4atnNfMHwVOiKOkzp1daG3Y+xTs6PFx9dueHvoXFQJQQ+3r0kYkLtsdYWsl5exIMvB0MWFHzP7LYjWhc7RN6YEA/ALI= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1701357315; c=relaxed/simple; bh=ZIDuJpfxHNqsItV9aHvuRs42ic4UPIfDJrBg3MAuj1o=; h=From:To:Subject:Date:Message-ID:MIME-Version; b=WeSyVCbI6ge7rI6D8om7nSRrD+VvDmmsVqdPs/Blz1/4YHUw//7bCBOoF+EkD9jKnNS5pBNyp2DOCh9erOdqTLXdErKtII5G4flZc9WjAcFcU1qLQ8JpniWaS9edl9zX8JVK9XlumoiZrIDikjgtUI/NuGTQmrYdh6iw5+noZuw= ARC-Authentication-Results: i=1; server2.sourceware.org X-CSE-ConnectionGUID: QT/QpX24Tc6UNj4y7woaqg== X-CSE-MsgGUID: k/L2fs1URW+tmIqBeMPM5w== X-IronPort-AV: E=Sophos;i="6.04,239,1695715200"; d="scan'208,223";a="27272311" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 30 Nov 2023 07:15:12 -0800 IronPort-SDR: Iap7xLI9Unjt1fITomOKdCxiF2DvNi5j6JXmRpx+5eCKnxCrjcBAwZHCF5RRcYSvtbKhDLzB5D catVbb2RdYzmgTzi0oHb7ZRdnhy4eXSt4AwroETS2ZYxbpJ2dJZ/HmUwylNHQZ7IGD3x3VLhYv Fpp+MfzTb80d2BtWSjaii4Z9wNq9xeEiC4UIe1uAoY1jK58ws8YhO28f+PJrZKHNlctn2CMdI4 g5jZYi3LhGlC3Oun9BVCAhvsUaoU2FKRc/CvsazQ+l+rBXl2Zzqh+8dq5QQre0fLOXgTCxB3yw dy8= From: Thomas Schwinge To: Paul-Antoine Arras , CC: Jakub Jelinek , Tobias Burnus Subject: Fix 'libgomp.c/declare-variant-4-*.c', add 'libgomp.c/declare-variant-4.c' (was: [PATCH] amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors) In-Reply-To: References: User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/27.1 (x86_64-pc-linux-gnu) Date: Thu, 30 Nov 2023 16:15:04 +0100 Message-ID: <8734wn2ryf.fsf@dem-tschwing-1.ger.mentorg.com> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) To svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) X-Spam-Status: No, score=-11.8 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.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.30 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 X-getmail-retrieved-from-mailbox: INBOX X-GMAIL-THRID: 1784002472670557420 X-GMAIL-MSGID: 1784002472670557420 Hi! On 2022-11-29T16:56:21+0100, Paul-Antoine Arras wrote: > [...] also adds test > cases checking all supported AMD ISAs are properly recognised when used > in a 'declare variant' construct. \o/ Yay for test cases! > --- /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" } } */ Etc. Unfortunately, I only ever see: UNSUPPORTED: libgomp.c/declare-variant-4-fiji.c UNSUPPORTED: libgomp.c/declare-variant-4-gfx803.c UNSUPPORTED: libgomp.c/declare-variant-4-gfx900.c UNSUPPORTED: libgomp.c/declare-variant-4-gfx906.c UNSUPPORTED: libgomp.c/declare-variant-4-gfx908.c UNSUPPORTED: libgomp.c/declare-variant-4-gfx90a.c Pushed to master branch commit aae57a9e19ba5d2016fa8a88db6d7ff15a40ca35 "Fix 'libgomp.c/declare-variant-4-*.c', add 'libgomp.c/declare-variant-4.c'", see attached. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 From aae57a9e19ba5d2016fa8a88db6d7ff15a40ca35 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Tue, 21 Nov 2023 16:36:48 +0100 Subject: [PATCH] Fix 'libgomp.c/declare-variant-4-*.c', add 'libgomp.c/declare-variant-4.c' These test cases being 'dg-skip-if [...] { ! amdgcn-*-* }' meant they were only ever considered for 'amdgcn-*-*' target -- that is, GCN *target*, not GCN *offload target*, what is intended to be tested here, but for which they thus always were UNSUPPORTED. Use the same style of 'dg-[...]' directives as for the nvptx offloading test cases, 'libgomp.c/declare-variant-3*.c'. Fix-up for commit 1fd508744eccda9ad9c6d6fcce5b2ea9c568818d "amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors". libgomp/ * testsuite/libgomp.c/declare-variant-4-fiji.c: Adjust. * testsuite/libgomp.c/declare-variant-4-gfx803.c: Likewise. * testsuite/libgomp.c/declare-variant-4-gfx900.c: Likewise. * testsuite/libgomp.c/declare-variant-4-gfx906.c: Likewise. * testsuite/libgomp.c/declare-variant-4-gfx908.c: Likewise. * testsuite/libgomp.c/declare-variant-4-gfx90a.c: Likewise. * testsuite/libgomp.c/declare-variant-4.h: Likewise. * testsuite/libgomp.c/declare-variant-4.c: New. --- libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c | 5 +++-- libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c | 5 +++-- libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c | 5 +++-- libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c | 5 +++-- libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c | 5 +++-- libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c | 5 +++-- libgomp/testsuite/libgomp.c/declare-variant-4.c | 8 ++++++++ libgomp/testsuite/libgomp.c/declare-variant-4.h | 5 +++++ 8 files changed, 31 insertions(+), 12 deletions(-) create mode 100644 libgomp/testsuite/libgomp.c/declare-variant-4.c diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c b/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c index 8a4e0f4728c..a138fb092f8 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c @@ -1,5 +1,6 @@ -/* { dg-do run { target { offload_target_amdgcn } } } */ -/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=fiji } */ /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ #define USE_FIJI_FOR_GFX803 diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c index 050d7c9dd79..03dffddac49 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c @@ -1,5 +1,6 @@ -/* { dg-do run { target { offload_target_amdgcn } } } */ -/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=fiji } */ /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ #include "declare-variant-4.h" diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c index 2eeb4a248c1..f3f5244c7bc 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c @@ -1,5 +1,6 @@ -/* { dg-do run { target { offload_target_amdgcn } } } */ -/* { dg-skip-if "gfx900 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx900" } } */ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx900 } */ /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ #include "declare-variant-4.h" diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c index 73a670dcc2a..ac43388a59f 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c @@ -1,5 +1,6 @@ -/* { dg-do run { target { offload_target_amdgcn } } } */ -/* { dg-skip-if "gfx906 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx906" } } */ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx906 } */ /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ #include "declare-variant-4.h" diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c index ead330f9f2c..f60741f202f 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c @@ -1,5 +1,6 @@ -/* { dg-do run { target { offload_target_amdgcn } } } */ -/* { dg-skip-if "gfx908 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx908" } } */ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx908 } */ /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ #include "declare-variant-4.h" diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c index a9b2d62a49d..832d174e4a5 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c @@ -1,5 +1,6 @@ -/* { dg-do run { target { offload_target_amdgcn } } } */ -/* { dg-skip-if "gfx90a only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx90a" } } */ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx90a } */ /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ #include "declare-variant-4.h" diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4.c b/libgomp/testsuite/libgomp.c/declare-variant-4.c new file mode 100644 index 00000000000..3c8802a6081 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4.c @@ -0,0 +1,8 @@ +/* { dg-additional-options -DOFFLOAD_DEVICE_GCN { target offload_device_gcn } } */ +/* { dg-additional-options {-fdump-tree-optimized -foffload-options=-fdump-tree-optimized} } */ + +#include "declare-variant-4.h" + +/* { dg-final { scan-tree-dump "= f \\(\\);" "optimized" } } + { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx\[^ \]+ \\(\\);" "optimized" { target offload_target_amdgcn } } } + { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "= f \\(\\);" "optimized" { target offload_target_nvptx } } } */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4.h b/libgomp/testsuite/libgomp.c/declare-variant-4.h index 47517b75ee7..a70352430c2 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4.h +++ b/libgomp/testsuite/libgomp.c/declare-variant-4.h @@ -61,10 +61,15 @@ main (void) #pragma omp target map(from : v) v = f (); +#ifdef OFFLOAD_DEVICE_GCN if (v == 0) __builtin_abort (); __builtin_printf ("AMDGCN accelerator: gfx%x\n", v); +#else + if (v != 0) + __builtin_abort (); +#endif return 0; } -- 2.34.1