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: 81038 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 9F830385AE4E for ; Thu, 30 Nov 2023 15:15:30 +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+patchwork=sourceware.org@gcc.gnu.org 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