aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorMartin Jambor <mjambor@suse.cz>2016-03-07 18:37:20 +0000
committerMartin Jambor <mjambor@suse.cz>2016-03-07 18:37:20 +0000
commitacda70b5427a872010a6048d1da6c780716909db (patch)
tree44108b98f60490d9cd387fe165fd23dc2b3298fc
parent6683af54ac10d086026040b1501437b4bfbe46e9 (diff)
[hsa testsuite] Gridification tests
2016-03-07 Martin Jambor <mjambor@suse.cz> * lib/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. git-svn-id: https://gcc.gnu.org/svn/gcc/trunk@234043 138bc75d-0d04-0410-961f-82ee72b054a4
-rw-r--r--gcc/testsuite/ChangeLog6
-rw-r--r--gcc/testsuite/c-c++-common/gomp/gridify-1.c54
-rw-r--r--gcc/testsuite/gfortran.dg/gomp/gridify-1.f9016
-rw-r--r--gcc/testsuite/lib/target-supports.exp8
4 files changed, 84 insertions, 0 deletions
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index f554c3a4188..faeff187e7d 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,9 @@
+2016-03-07 Martin Jambor <mjambor@suse.cz>
+
+ * lib/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.
+
2016-03-07 Andre Vieira <andre.simoesdiasvieira@arm.com>
* gcc.target/arm/pr45701-1.c: Change assembler scan to not
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 00000000000..ba7a86665b5
--- /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 00000000000..00ff7f510a0
--- /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 14a8bc8d88c..16d358882d5 100644
--- a/gcc/testsuite/lib/target-supports.exp
+++ b/gcc/testsuite/lib/target-supports.exp
@@ -6937,3 +6937,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" ]
+}