[hsa,testsuite] Introduce offload_device_shared_as effective target

2016-02-26  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.

Co-Authored-By: Martin Jambor <mjambor@suse.cz>
From-SVN: r233757

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 3cf28601f0d7a79d0c4bf8335c429bdcf6753f43..e96f1495fc255044b19c70e65b2952fcef6250b8 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,10 @@
+2016-02-26  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.
+
 2016-02-25  Ilya Verbin  <ilya.verbin@intel.com>
 
        PR driver/68463

diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index a4c9d834e461aa96cebb15ac90e07668d98da297..154a44769542d687919423ee0589fbbbf52fc0b8 100644 (file)
--- 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 (file)
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;
+}