[hsa, testsuite] Gridification tests

[Martin Jambor <mjambor at suse dot cz>]

Hi,

the patch below adds a DejaGNU effective target predicate (is that the
correct dejagnu term?) offload_hsa so that selected tests can be run
only if the hsa offloading is enabled.  I hope it is fairly standard
stuff.  Additionally, it adds one C/C++ and one Fortran testsuite to
check that gridification happens.

Tested, both with and without HSA enabled.  OK for trunk?

Thanks,

Martin

2016-02-10  Martin Jambor  <mjambor@suse.cz>

	* target-supports.exp (check_effective_target_offload_hsa): New.
	* c-c++-common/gomp/gridify-1.c: New test.
        * gfortran.dg/gomp/gridify-1.f90: Likewise.
---
 gcc/testsuite/c-c++-common/gomp/gridify-1.c  | 54 ++++++++++++++++++++++++++++
 gcc/testsuite/gfortran.dg/gomp/gridify-1.f90 | 16 +++++++++
 gcc/testsuite/lib/target-supports.exp        |  8 +++++
 3 files changed, 78 insertions(+)
 create mode 100644 gcc/testsuite/c-c++-common/gomp/gridify-1.c
 create mode 100644 gcc/testsuite/gfortran.dg/gomp/gridify-1.f90

diff --git a/gcc/testsuite/c-c++-common/gomp/gridify-1.c b/gcc/testsuite/c-c++-common/gomp/gridify-1.c
new file mode 100644
index 0000000..ba7a866
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/gridify-1.c
@@ -0,0 +1,54 @@
+/* { dg-do compile } */
+/* { dg-require-effective-target offload_hsa } */
+/* { dg-options "-fopenmp -fdump-tree-omplower-details" } */
+
+void
+foo1 (int n, int *a, int workgroup_size)
+{
+  int i;
+#pragma omp target
+#pragma omp teams thread_limit(workgroup_size)
+#pragma omp distribute parallel for shared(a) firstprivate(n) private(i)
+    for (i = 0; i < n; i++)
+      a[i]++;
+}
+
+void
+foo2 (int j, int n, int *a)
+{
+  int i;
+#pragma omp target teams
+#pragma omp distribute parallel for shared(a) firstprivate(n) private(i) firstprivate(j)
+    for (i = j + 1; i < n; i++)
+      a[i] = i;
+}
+
+void
+foo3 (int j, int n, int *a)
+{
+  int i;
+#pragma omp target teams
+#pragma omp distribute parallel for shared(a) firstprivate(n) private(i) firstprivate(j)
+  for (i = j + 1; i < n; i += 3)
+    a[i] = i;
+}
+
+void
+foo4 (int j, int n, int *a)
+{
+#pragma omp parallel
+  {
+    #pragma omp single
+    {
+      int i;
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for shared(a) firstprivate(n) private(i) firstprivate(j)
+      for (i = j + 1; i < n; i += 3)
+	a[i] = i;
+    }
+  }
+}
+
+
+/* { dg-final { scan-tree-dump-times "Target construct will be turned into a gridified GPGPU kernel" 4 "omplower" } } */
diff --git a/gcc/testsuite/gfortran.dg/gomp/gridify-1.f90 b/gcc/testsuite/gfortran.dg/gomp/gridify-1.f90
new file mode 100644
index 0000000..00ff7f5
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/gridify-1.f90
@@ -0,0 +1,16 @@
+! { dg-do compile }
+! { dg-require-effective-target offload_hsa }
+! { dg-options "-fopenmp -fdump-tree-omplower-details" } */
+
+subroutine vector_square(n, a, b)
+      integer i, n, b(n), a(n)
+!$omp target teams
+!$omp distribute parallel do
+      do i=1,n
+          b(i) = a(i) * a(i)
+      enddo
+!$omp end distribute parallel do
+!$omp end target teams
+end subroutine vector_square
+
+! { dg-final { scan-tree-dump "Target construct will be turned into a gridified GPGPU kernel" "omplower" } }
diff --git a/gcc/testsuite/lib/target-supports.exp b/gcc/testsuite/lib/target-supports.exp
index 0b4252f..fac4c3c 100644
--- a/gcc/testsuite/lib/target-supports.exp
+++ b/gcc/testsuite/lib/target-supports.exp
@@ -6936,3 +6936,11 @@ proc check_effective_target_offload_nvptx { } {
 	int main () {return 0;}
     } "-foffload=nvptx-none" ]
 }
+
+# Return 1 if the compiler has been configured with hsa offloading.
+
+proc check_effective_target_offload_hsa { } {
+    return [check_no_compiler_messages offload_hsa assembly {
+	int main () {return 0;}
+    } "-foffload=hsa" ]
+}
-- 
2.7.1