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: 1957 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a5d:4ac7:0:0:0:0:0 with SMTP id y7csp2630905wrs; Wed, 12 Oct 2022 07:06:33 -0700 (PDT) X-Google-Smtp-Source: AMsMyM4g5kt3A9Kj9oVneuIbqiu4+iuM+yjN07EAh+yHW4pxTh31yGBRA5YvBMl/XQl6hzXFAYp+ X-Received: by 2002:a05:6402:4303:b0:45c:cd3d:f5 with SMTP id m3-20020a056402430300b0045ccd3d00f5mr668136edc.188.1665583593564; Wed, 12 Oct 2022 07:06:33 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1665583593; cv=none; d=google.com; s=arc-20160816; b=IqZ/YisaNjesMCnIElbZyZiwiHzfDcFtHFTYSamNRtHf6NQ4SXhMF9jKdJrjZeSTPk jX7Rgy18EckiiJWBevQ/CY9Dhs5xy1uomIPcp82NyMA3XbYkZBLcdf3qRTzzg9Eq+TOX ge3SIgsiP+olWHBxetzNjP6Hir2fvounokGNagasvO37Y144inPPRXZVxmM5lVPuKBAb 2+vkWgeh3xGdfo9vcuKyHHW+ZQ6tj6LtGNAMYkfHvWTQp9Vtun14/jTZlhKFmh6TPBgy r1HknXkDoszsd7KrqluumjxnR0MTXXlPjsXvRDOBo64KO2iiv4rczDJ/M5dyHQKMtGsr aYFQ== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:cc: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=aArBIx4xbef+/yBwYUvj0UvjK479VKCS7iFwgnqOVAs=; b=vGl/YHXLeRiVgpJVDff2DoXHswbal1A6sRuGLtdlkYGjS6yurArQQpbXIaEmtlJQ6T 8KWg87d1u4oj3L6AjWZawqzrq6CJXVU4lH0cRrkqUOpcH4JCDAkUnED8jMiF8Y/Bul95 kBhVl1lHDMJm9W4o4Xi+sKj7i0LtI0Q3jXeB/cbqMH/Y+uiqOrMR0i2tJms17Hf8kAKX JzTmsgE4jLccxB4WJtj6WOeJfQykeW3QJSGuI9o5sqKQR26m4dESDSO1fSVErJEoxCOv i/sYaYfu4irr1/zPC9vmoBrixBETdGj1BUqNjl7emRiZU2rUI2dWQh7wt4tSJ/UInf9y Bg1Q== 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 (ip-8-43-85-97.sourceware.org. [8.43.85.97]) by mx.google.com with ESMTPS id hq13-20020a1709073f0d00b0078d20d71475si3170384ejc.413.2022.10.12.07.06.33 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 12 Oct 2022 07:06:33 -0700 (PDT) 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 8F9353850216 for ; Wed, 12 Oct 2022 14:06:28 +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+ouuuleilei=gmail.com@gcc.gnu.org Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: =?utf-8?q?INBOX?= X-GMAIL-THRID: =?utf-8?q?1746490982246852848?= X-GMAIL-MSGID: =?utf-8?q?1746490982246852848?= 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; +}