From 622174b20dfaa584abcf950c89aaa1843b146669 Mon Sep 17 00:00:00 2001 From: Keith McDaniel Date: Fri, 26 Feb 2016 19:42:31 +0000 Subject: [PATCH] [hsa,testsuite] Introduce offload_device_shared_as effective target 2016-02-26 Keith McDaniel Martin Jambor * testsuite/lib/libgomp.exp (check_effective_target_offload_device_shared_as): New proc. * testsuite/libgomp.c++/declare_target-1.C: New test. Co-Authored-By: Martin Jambor From-SVN: r233757 --- libgomp/ChangeLog | 7 ++++ libgomp/testsuite/lib/libgomp.exp | 13 +++++++ .../testsuite/libgomp.c++/declare_target-1.C | 38 +++++++++++++++++++ 3 files changed, 58 insertions(+) create mode 100644 libgomp/testsuite/libgomp.c++/declare_target-1.C diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 3cf28601f0d7..e96f1495fc25 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,10 @@ +2016-02-26 Keith McDaniel + Martin Jambor + + * testsuite/lib/libgomp.exp + (check_effective_target_offload_device_shared_as): New proc. + * testsuite/libgomp.c++/declare_target-1.C: New test. + 2016-02-25 Ilya Verbin PR driver/68463 diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index a4c9d834e461..154a44769542 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 000000000000..4394bb1405df --- /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 + +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.47.2