From patchwork Wed Oct 12 14:05:32 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 58701 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 3E0283821FF3 for ; Wed, 12 Oct 2022 14:06:21 +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 B445438582BC for ; Wed, 12 Oct 2022 14:06:01 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org B445438582BC 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.95,179,1661846400"; d="diff'?scan'208";a="84660686" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 12 Oct 2022 06:05:40 -0800 IronPort-SDR: +L0WrLzx01grppXstU9v2MPacXr3UO9UTTipOu+vvoBcg6iulrs6NRii7WfEfnr42poN7dlASt PVfc7DHV9eeLxIGlY5KKIoXf8+10JuS8cczk3oPG85letv9qZel0GYj2fvetN5KLsXw+31P1pb Blhb55GlWIJDG+c3vKnFHnAGagI//cF3KTOrNOM3GqXKxtCVbtPi7tIo8IwR5/e6KTYO9V9REw T1TJalGFYMnrsOl5ht0lGZLEmmh/+QNm6fruA8ye0g2ZpBGoom35GFoNd63yaT1S0dH0HkZFP+ PNE= Message-ID: Date: Wed, 12 Oct 2022 16:05:32 +0200 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:102.0) Gecko/20100101 Thunderbird/102.3.2 Content-Language: en-US To: gcc-patches , Jakub Jelinek , Andrew Stubbs From: Tobias Burnus Subject: [Patch] libgomp: Add offload_device_gcn check, add requires-4a.c test X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-08.mgc.mentorg.com (139.181.222.8) To svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) X-Spam-Status: No, score=-11.4 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: , Cc: "Vollweiler, Marcel" Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" This came up because the USM implementation with -foffload-memory={unified,pinned} as posted at https://gcc.gnu.org/pipermail/gcc-patches/2022-July/597976.html does not handle USM with static variables. This shows up for the OG12 alias devel/omp/gcc-12 branch as FAIL for requires-4.c. The attached patch prepares for skipping requires-4.c for the gcn/nvptx device and adds an adjacent requires-4a.c testcase, using heap memory, that can still run on gcn/nvptx. Additionally, I commented on no longer used #defined, following the precedence GOMP_DEVICE_HOST_NONSHM. Thus, this tests adds another testcase and one effective-target check, out-comments a unused #define - and that's it. (Otherwise, it is just a prep patch.) OK for mainline? Tobias PS: Currently, neither the preexisting offload_device_nvptx nor the new offload_device_gcn target selector is used, neither in old code nor by this patch. ----------------- 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 libgomp: Add offload_device_gcn check, add requires-4a.c test Duplicate libgomp.c-c++-common/requires-4.c (as ...-4a.c) but with using a heap-allocated instead of static memory for a variable. This change and the added offload_device_gcn check prepare for pseudo-USM, where the device hardware cannot access all host memory but only managed and pinned memory; for those, requires-4.c will fail and the new check permits to add target { ! { offload_device_nvptx || offload_device_gcn } } to requires-4.c; however, it has not been added yet as pseuo-USM support is not yet on mainline. (Review is pending for the USM patches.) include/ChangeLog: * gomp-constants.h (GOMP_DEVICE_HSA): Comment (unused). libgomp/ChangeLog: * testsuite/lib/libgomp.exp (check_effective_target_offload_device_gcn): New. * testsuite/libgomp.c-c++-common/on_device_arch.h (device_arch_gcn, on_device_arch_gcn): New. * testsuite/libgomp.c-c++-common/requires-4a.c: New test; copied from requires-4.c but using heap-allocated memory. include/gomp-constants.h | 2 +- libgomp/testsuite/lib/libgomp.exp | 12 +++++++ .../libgomp.c-c++-common/on_device_arch.h | 13 ++++++++ .../testsuite/libgomp.c-c++-common/requires-4a.c | 39 ++++++++++++++++++++++ 4 files changed, 65 insertions(+), 1 deletion(-) diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 84316f953d0..fac7316b858 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -229,9 +229,9 @@ enum gomp_map_kind /* #define GOMP_DEVICE_HOST_NONSHM 3 removed. */ #define GOMP_DEVICE_NOT_HOST 4 #define GOMP_DEVICE_NVIDIA_PTX 5 #define GOMP_DEVICE_INTEL_MIC 6 -#define GOMP_DEVICE_HSA 7 +/* #define GOMP_DEVICE_HSA 7 removed. */ #define GOMP_DEVICE_GCN 8 /* We have a compatibility issue. OpenMP 5.2 introduced omp_initial_device with value of -1 which clashes with our diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index 107a3c2ac9d..4b8c64de8a5 100644 --- a/libgomp/testsuite/lib/libgomp.exp +++ b/libgomp/testsuite/lib/libgomp.exp @@ -414,8 +414,20 @@ proc check_effective_target_offload_device_nvptx { } { } } ] } +# Return 1 if using a GCN offload device. +proc check_effective_target_offload_device_gcn { } { + return [check_runtime_nocache offload_device_gcn { + #include + #include "testsuite/libgomp.c-c++-common/on_device_arch.h" + int main () + { + return !on_device_arch_gcn (); + } + } ] +} + # Return 1 if at least one Nvidia GPU is accessible. proc check_effective_target_openacc_nvidia_accel_present { } { return [check_runtime openacc_nvidia_accel_present { diff --git a/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h b/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h index f92743b04d7..6f66dbd784c 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h +++ b/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h @@ -6,15 +6,22 @@ device_arch_nvptx (void) { return GOMP_DEVICE_NVIDIA_PTX; } +/* static */ int +device_arch_gcn (void) +{ + return GOMP_DEVICE_GCN; +} + /* static */ int device_arch_intel_mic (void) { return GOMP_DEVICE_INTEL_MIC; } #pragma omp declare variant (device_arch_nvptx) match(construct={target},device={arch(nvptx)}) +#pragma omp declare variant (device_arch_gcn) match(construct={target},device={arch(gcn)}) #pragma omp declare variant (device_arch_intel_mic) match(construct={target},device={arch(intel_mic)}) /* static */ int device_arch (void) { @@ -36,8 +43,14 @@ on_device_arch_nvptx () { return on_device_arch (GOMP_DEVICE_NVIDIA_PTX); } +int +on_device_arch_gcn () +{ + return on_device_arch (GOMP_DEVICE_GCN); +} + int on_device_arch_intel_mic () { return on_device_arch (GOMP_DEVICE_INTEL_MIC); diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-4a.c b/libgomp/testsuite/libgomp.c-c++-common/requires-4a.c new file mode 100644 index 00000000000..4fb9783a97a --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-4a.c @@ -0,0 +1,39 @@ +/* { dg-additional-options "-flto" } */ +/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */ +/* { dg-additional-sources requires-4-aux.c } */ + +/* Same as requires-4.c, but uses heap memory for 'a'. */ + +/* Check no diagnostic by device-compiler's or host compiler's lto1. + Other file uses: 'requires reverse_offload', but that's inactive as + there are no declare target directives, device constructs nor device routines */ + +/* Depending on offload device capabilities, it may print something like the + following (only) if GOMP_DEBUG=1: + "devices present but 'omp requires unified_address, unified_shared_memory, reverse_offload' cannot be fulfilled" + and in that case does host-fallback execution. + + No offload devices support USM at present, so we may verify host-fallback + execution by presence of separate memory spaces. */ + +#pragma omp requires unified_address,unified_shared_memory + +int *a; +extern void foo (void); + +int +main (void) +{ + a = (int *) __builtin_calloc (sizeof (int), 10); + #pragma omp target map(to: a) + for (int i = 0; i < 10; i++) + a[i] = i; + + for (int i = 0; i < 10; i++) + if (a[i] != i) + __builtin_abort (); + + foo (); + __builtin_free (a); + return 0; +}