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