Hello, this patch has been written by Keith McDaniel when he was working at AMD (and so it should be covered by their blanket copyright assignment) and adds a libgomp testsuite effective target predicate offload_device_shared_as that allows us to run tests only when target constructs are run on a device with shared memory (this includes the host). Keith included a C++ test to illustrate the use.
I have tested this thoroughly, both with and without HSA (enabled or present). OK for trunk? Thanks, Martin 2016-02-10 Keith McDaniel <k.allen.mcdan...@gmail.com> Martin Jambor <mjam...@suse.cz> * testsuite/lib/libgomp.exp (check_effective_target_offload_device_shared_as): New proc. * testsuite/libgomp.c++/declare_target-1.C: New test. --- libgomp/testsuite/lib/libgomp.exp | 13 ++++++++ libgomp/testsuite/libgomp.c++/declare_target-1.C | 38 ++++++++++++++++++++++++ 2 files changed, 51 insertions(+) create mode 100644 libgomp/testsuite/libgomp.c++/declare_target-1.C diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index a4c9d83..154a447 100644 --- a/libgomp/testsuite/lib/libgomp.exp +++ b/libgomp/testsuite/lib/libgomp.exp @@ -343,6 +343,19 @@ proc check_effective_target_offload_device_nonshared_as { } { } } ] } + +# Return 1 if offload device is available and it has shared address space. +proc check_effective_target_offload_device_shared_as { } { + return [check_runtime_nocache offload_device_shared_as { + int main () + { + int x = 10; + #pragma omp target map(to: x) + x++; + return x == 10; + } + } ] +} # Return 1 if at least one nvidia board is present. diff --git a/libgomp/testsuite/libgomp.c++/declare_target-1.C b/libgomp/testsuite/libgomp.c++/declare_target-1.C new file mode 100644 index 0000000..4394bb1 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare_target-1.C @@ -0,0 +1,38 @@ +// { dg-do run } +// { dg-require-effective-target offload_device_shared_as } + +#include <stdlib.h> + +struct typeX +{ + int a; +}; + +class typeY +{ +public: + int foo () { return a^0x01; } + int a; +}; + +#pragma omp declare target +struct typeX varX; +class typeY varY; +#pragma omp end declare target + +int main () +{ + varX.a = 0; + varY.a = 0; + + #pragma omp target + { + varX.a = 100; + varY.a = 100; + } + + if (varX.a != 100 || varY.a != 100) + abort (); + + return 0; +} -- 2.7.1