[hsa,testsuite] Introduce offload_device_shared_as effective target

[Martin Jambor <mjambor at suse dot cz>]

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.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.
---
 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