diff options
author | Tobias Burnus <tobias@codesourcery.com> | 2022-10-20 12:58:52 +0200 |
---|---|---|
committer | Tobias Burnus <tobias@codesourcery.com> | 2022-10-20 12:58:52 +0200 |
commit | 12d9f5afbd2660862045acd41cb65a77e35bea4d (patch) | |
tree | 9f6ce9c506c62938f9f40147dac4fa259e6fe7af /libgomp | |
parent | 5362b5cc8dba9de70904613fff1d5391a9c087ce (diff) | |
download | gcc-12d9f5afbd2660862045acd41cb65a77e35bea4d.zip gcc-12d9f5afbd2660862045acd41cb65a77e35bea4d.tar.gz gcc-12d9f5afbd2660862045acd41cb65a77e35bea4d.tar.bz2 |
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.
Diffstat (limited to 'libgomp')
-rw-r--r-- | libgomp/testsuite/lib/libgomp.exp | 12 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h | 13 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.c-c++-common/requires-4a.c | 39 |
3 files changed, 64 insertions, 0 deletions
diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index 107a3c2..4b8c64d 100644 --- a/libgomp/testsuite/lib/libgomp.exp +++ b/libgomp/testsuite/lib/libgomp.exp @@ -415,6 +415,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 <omp.h> + #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 f92743b..6f66dbd 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h +++ b/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h @@ -8,12 +8,19 @@ device_arch_nvptx (void) } /* 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) @@ -38,6 +45,12 @@ on_device_arch_nvptx () } 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 0000000..4fb9783 --- /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; +} |