Message ID | 20160226155941.GI3094@virgil.suse.cz |
---|---|
State | New |
Headers | show |
On Fri, Feb 26, 2016 at 04:59:42PM +0100, Martin Jambor wrote: > 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.mcdaniel@gmail.com> > Martin Jambor <mjambor@suse.cz> > > * testsuite/lib/libgomp.exp > (check_effective_target_offload_device_shared_as): New proc. > * testsuite/libgomp.c++/declare_target-1.C: New test. Ok. Jakub
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; +}