diff options
author | Martin Jambor <mjambor@suse.cz> | 2016-03-07 18:37:20 +0000 |
---|---|---|
committer | Martin Jambor <mjambor@suse.cz> | 2016-03-07 18:37:20 +0000 |
commit | acda70b5427a872010a6048d1da6c780716909db (patch) | |
tree | 44108b98f60490d9cd387fe165fd23dc2b3298fc | |
parent | 6683af54ac10d086026040b1501437b4bfbe46e9 (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/ChangeLog | 6 | ||||
-rw-r--r-- | gcc/testsuite/c-c++-common/gomp/gridify-1.c | 54 | ||||
-rw-r--r-- | gcc/testsuite/gfortran.dg/gomp/gridify-1.f90 | 16 | ||||
-rw-r--r-- | gcc/testsuite/lib/target-supports.exp | 8 |
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" ] +} |