From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1534) id 3C5D138561BE; Thu, 20 Oct 2022 11:08:10 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 3C5D138561BE DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1666264090; bh=gFK0T3X9Km2VTxcfuD0h0AD94j/ZXkGPdwXG85FUd98=; h=From:To:Subject:Date:From; b=Px0Z1HM4xvY7dlgCmeDQvmRkoq/JVhmE570CQR25LiivBgttvPK/h5MzCmhkstB52 03czjkU6qIiZP61kWAmhpoEhZ0VmZb964+RK7NQl3jxJdsFhW/Oh9FFiU4KpbqsuV+ JStX7dmEdBo3u8x5tRLc0yEjFjgt9pFiK6cM1bvU= Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Tobias Burnus To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-12] libgomp: Add offload_device_gcn check, add requires-4a.c test X-Act-Checkin: gcc X-Git-Author: Tobias Burnus X-Git-Refname: refs/heads/devel/omp/gcc-12 X-Git-Oldrev: 0d6fc5032c7ba8a95301d0ccbc418875e73955ac X-Git-Newrev: ffc6267a829068f49fd4346a0495d3f3fc51c89e Message-Id: <20221020110810.3C5D138561BE@sourceware.org> Date: Thu, 20 Oct 2022 11:08:10 +0000 (GMT) List-Id: https://gcc.gnu.org/g:ffc6267a829068f49fd4346a0495d3f3fc51c89e commit ffc6267a829068f49fd4346a0495d3f3fc51c89e Author: Tobias Burnus Date: Thu Oct 20 13:07:37 2022 +0200 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 out unused define. 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. (cherry picked from commit 12d9f5afbd2660862045acd41cb65a77e35bea4d) Diff: --- include/ChangeLog.omp | 7 ++++ include/gomp-constants.h | 2 +- libgomp/ChangeLog.omp | 12 +++++++ libgomp/testsuite/lib/libgomp.exp | 12 +++++++ .../libgomp.c-c++-common/on_device_arch.h | 13 ++++++++ .../testsuite/libgomp.c-c++-common/requires-4a.c | 39 ++++++++++++++++++++++ 6 files changed, 84 insertions(+), 1 deletion(-) diff --git a/include/ChangeLog.omp b/include/ChangeLog.omp index 141ac55cf7a..4a72c8b9af5 100644 --- a/include/ChangeLog.omp +++ b/include/ChangeLog.omp @@ -1,3 +1,10 @@ +2022-10-20 Tobias Burnus + + Backport from mainline: + 2022-10-20 Tobias Burnus + + * gomp-constants.h (GOMP_DEVICE_HSA): Comment out unused define. + 2022-08-31 Tobias Burnus Revert: diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 918414ec218..dd753a82106 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -257,7 +257,7 @@ enum gomp_map_kind #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 diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index 570aabf82c4..cb884aa0c83 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,15 @@ +2022-10-20 Tobias Burnus + + Backport from mainline: + 2022-10-20 Tobias Burnus + + * 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. + 2022-10-17 Tobias Burnus Backport from mainline: diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index c98c1edf57b..ffdcfd85b72 100644 --- a/libgomp/testsuite/lib/libgomp.exp +++ b/libgomp/testsuite/lib/libgomp.exp @@ -414,6 +414,18 @@ 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 { } { 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 @@ -7,6 +7,12 @@ 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) { @@ -14,6 +20,7 @@ device_arch_intel_mic (void) } #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) @@ -37,6 +44,12 @@ 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 () { 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; +}