[hsa testsuite 1/5] Gridification tests

[hsa testsuite 1/5] Gridification tests

[Martin Jambor]

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

[hsa testsuite 3/5] Suppress hsa warnings in libgomp tests

[Martin Jambor]

Hi,

just like with the compiler gomp testsuite, we need to add -Wno-hsa to
options when compiling libgomp testcases in order not to have "excess
errors" failures when HSA is enabled.  There are quite many of such
testcases on the trunk because I have disabled the dynamic parallelism
way of executing stuff.

The patch below adds the option to all libgomp testsuite compilations,
so that people who are not interested in HSA do not need to care.  The
patch has been tested both with and without HSA enabled.  OK for
trunk?

Thanks,

Martin


2016-03-04  Martin Jambor  <mjambor@suse.cz>

	* testsuite/lib/libgomp.exp (libgomp_init): Append -Wno-hsa to
	ALWAYS_CFLAGS.
---
 libgomp/testsuite/lib/libgomp.exp | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index 154a447..bbc2c26 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -237,6 +237,9 @@ proc libgomp_init { args } {
     # Disable caret
     lappend ALWAYS_CFLAGS "additional_flags=-fno-diagnostics-show-caret"
 
+    # Disable HSA warnings by default.
+    lappend ALWAYS_CFLAGS "additional_flags=-Wno-hsa"
+
     # Disable color diagnostics
     lappend ALWAYS_CFLAGS "additional_flags=-fdiagnostics-color=never"
 
-- 
2.7.1

[hsa testsuite 4/5] Adjust libgomp tests that do not work on host fallback

[Martin Jambor]

Hi,

this patch avoids run-time failures in libgomp testsuite that
curtrently happen when HSA offloading is actually used.  All of these
tests require the offload_device effective target which the patch
changes to offload_device_nonshared_as one.

For some tests, such as libgomp.c/examples-4/device-1.c this is
clearly just the correct thing to do because the test explicitely
checks that changes that happen in a target construct and are not
"mapped" back are not observable on the host.

However, the majority of the tests has a different problem.  If a test
for some reason is not compiled into HSAIL (usually because it would
require the dynamic parallelism path which is disabled or because it
calls abort from within target which HSA so far cannot handle), the
host fallback is called, even though the test actually is not supposed
to be called on it.  Such problematic tests then call
omp_is_initial_device to verify they are not running on the host and
decide to fail when they figure out they are.

Changing the effective target only to devices with non-shared memory
probably isn't the really correct fix.  We basically want to disable
the host fallback for them regardeless of address spaces but I cannot
think of a simple and generic way of doing that.  However, all
testcases for non-shared memory devices were written with disallowed
fallback in mind and so this soulution also gives the desired result.
Perhaps we need something better for the long term, any suggestions
are welcome.

Tested both with and without HSA (enabled or present).  OK for trunk?

Thanks,

Martin

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

libgomp/
	* testsuite/libgomp.c/examples-4/async_target-2.c: Only run on
	non-shared memory accelerators.
	* testsuite/libgomp.c/examples-4/device-1.c: Likewise.
	* testsuite/libgomp.c/examples-4/target-5.c: Likewise.
	* testsuite/libgomp.c/examples-4/target_data-6.c: Likewise.
	* testsuite/libgomp.c/examples-4/target_data-7.c: Likewise.
	* testsuite/libgomp.fortran/examples-4/async_target-2.f90: Likewise.
	* testsuite/libgomp.fortran/examples-4/device-1.f90: Likewise.
	* testsuite/libgomp.fortran/examples-4/target-5.f90: Likewise.
	* testsuite/libgomp.fortran/examples-4/target_data-6.f90: Likewise.
	* testsuite/libgomp.fortran/examples-4/target_data-7.f90: Likewise.
---
 libgomp/testsuite/libgomp.c/examples-4/async_target-2.c         | 2 +-
 libgomp/testsuite/libgomp.c/examples-4/device-1.c               | 2 +-
 libgomp/testsuite/libgomp.c/examples-4/target-5.c               | 2 +-
 libgomp/testsuite/libgomp.c/examples-4/target_data-6.c          | 2 +-
 libgomp/testsuite/libgomp.c/examples-4/target_data-7.c          | 2 +-
 libgomp/testsuite/libgomp.fortran/examples-4/async_target-2.f90 | 2 +-
 libgomp/testsuite/libgomp.fortran/examples-4/device-1.f90       | 2 +-
 libgomp/testsuite/libgomp.fortran/examples-4/target-5.f90       | 2 +-
 libgomp/testsuite/libgomp.fortran/examples-4/target_data-6.f90  | 2 +-
 libgomp/testsuite/libgomp.fortran/examples-4/target_data-7.f90  | 2 +-
 10 files changed, 10 insertions(+), 10 deletions(-)

diff --git a/libgomp/testsuite/libgomp.c/examples-4/async_target-2.c b/libgomp/testsuite/libgomp.c/examples-4/async_target-2.c
index ce63328..0c76f8e 100644
--- a/libgomp/testsuite/libgomp.c/examples-4/async_target-2.c
+++ b/libgomp/testsuite/libgomp.c/examples-4/async_target-2.c
@@ -1,5 +1,5 @@
 /* { dg-do run } */
-/* { dg-require-effective-target offload_device } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
 
 #include <omp.h>
 #include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.c/examples-4/device-1.c b/libgomp/testsuite/libgomp.c/examples-4/device-1.c
index dad8572..46aa160 100644
--- a/libgomp/testsuite/libgomp.c/examples-4/device-1.c
+++ b/libgomp/testsuite/libgomp.c/examples-4/device-1.c
@@ -1,5 +1,5 @@
 /* { dg-do run } */
-/* { dg-require-effective-target offload_device } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
 
 #include <omp.h>
 #include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.c/examples-4/target-5.c b/libgomp/testsuite/libgomp.c/examples-4/target-5.c
index 1853fba..1c14bae 100644
--- a/libgomp/testsuite/libgomp.c/examples-4/target-5.c
+++ b/libgomp/testsuite/libgomp.c/examples-4/target-5.c
@@ -1,5 +1,5 @@
 /* { dg-do run } */
-/* { dg-require-effective-target offload_device } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
 
 #include <omp.h>
 #include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.c/examples-4/target_data-6.c b/libgomp/testsuite/libgomp.c/examples-4/target_data-6.c
index affeb49..57c7c0c 100644
--- a/libgomp/testsuite/libgomp.c/examples-4/target_data-6.c
+++ b/libgomp/testsuite/libgomp.c/examples-4/target_data-6.c
@@ -1,5 +1,5 @@
 /* { dg-do run } */
-/* { dg-require-effective-target offload_device } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
 
 #include <stdlib.h>
 #include <omp.h>
diff --git a/libgomp/testsuite/libgomp.c/examples-4/target_data-7.c b/libgomp/testsuite/libgomp.c/examples-4/target_data-7.c
index c18d480..8ec41ea 100644
--- a/libgomp/testsuite/libgomp.c/examples-4/target_data-7.c
+++ b/libgomp/testsuite/libgomp.c/examples-4/target_data-7.c
@@ -1,5 +1,5 @@
 /* { dg-do run } */
-/* { dg-require-effective-target offload_device } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
 
 #include <stdlib.h>
 #include <omp.h>
diff --git a/libgomp/testsuite/libgomp.fortran/examples-4/async_target-2.f90 b/libgomp/testsuite/libgomp.fortran/examples-4/async_target-2.f90
index b12b0ea..a94db89 100644
--- a/libgomp/testsuite/libgomp.fortran/examples-4/async_target-2.f90
+++ b/libgomp/testsuite/libgomp.fortran/examples-4/async_target-2.f90
@@ -1,5 +1,5 @@
 ! { dg-do run }
-! { dg-require-effective-target offload_device }
+! { dg-require-effective-target offload_device_nonshared_as }
 
 subroutine init (v1, v2, N)
   !$omp declare target
diff --git a/libgomp/testsuite/libgomp.fortran/examples-4/device-1.f90 b/libgomp/testsuite/libgomp.fortran/examples-4/device-1.f90
index 291604b..a411db4 100644
--- a/libgomp/testsuite/libgomp.fortran/examples-4/device-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/examples-4/device-1.f90
@@ -1,5 +1,5 @@
 ! { dg-do run }
-! { dg-require-effective-target offload_device }
+! { dg-require-effective-target offload_device_nonshared_as }
 
 program e_57_1
   use omp_lib, only: omp_is_initial_device
diff --git a/libgomp/testsuite/libgomp.fortran/examples-4/target-5.f90 b/libgomp/testsuite/libgomp.fortran/examples-4/target-5.f90
index 3f454d7..813a273 100644
--- a/libgomp/testsuite/libgomp.fortran/examples-4/target-5.f90
+++ b/libgomp/testsuite/libgomp.fortran/examples-4/target-5.f90
@@ -1,5 +1,5 @@
 ! { dg-do run }
-! { dg-require-effective-target offload_device }
+! { dg-require-effective-target offload_device_nonshared_as }
 
 module e_50_5_mod
 integer, parameter :: THRESHOLD1 = 500, THRESHOLD2 = 100
diff --git a/libgomp/testsuite/libgomp.fortran/examples-4/target_data-6.f90 b/libgomp/testsuite/libgomp.fortran/examples-4/target_data-6.f90
index 258da21e..9b79104 100644
--- a/libgomp/testsuite/libgomp.fortran/examples-4/target_data-6.f90
+++ b/libgomp/testsuite/libgomp.fortran/examples-4/target_data-6.f90
@@ -1,5 +1,5 @@
 ! { dg-do run }
-! { dg-require-effective-target offload_device }
+! { dg-require-effective-target offload_device_nonshared_as }
 
 module e_51_6_mod
 integer, parameter :: THRESHOLD = 500
diff --git a/libgomp/testsuite/libgomp.fortran/examples-4/target_data-7.f90 b/libgomp/testsuite/libgomp.fortran/examples-4/target_data-7.f90
index 2ddac9e..8fc5832 100644
--- a/libgomp/testsuite/libgomp.fortran/examples-4/target_data-7.f90
+++ b/libgomp/testsuite/libgomp.fortran/examples-4/target_data-7.f90
@@ -1,5 +1,5 @@
 ! { dg-do run }
-! { dg-require-effective-target offload_device }
+! { dg-require-effective-target offload_device_nonshared_as }
 
 module e_51_7_mod
 integer, parameter :: THRESHOLD = 500
-- 
2.7.1

[hsa testsuite 5/5] New directory for HSA-specific C testcases

[Martin Jambor]

Hi,

we would like a place to have some HSA-specific tests, which would
only run not only when HSA is enabled at configuration time but also
when HSA hardware is present and used for offloading.

I have proposed the first version of this patch as
https://gcc.gnu.org/ml/gcc-patches/2016-02/msg01817.html and got some
seedback from Mike Stump in
https://gcc.gnu.org/ml/gcc-patches/2016-03/msg00086.html.  I hope I
have incorporated his suggestions.  As I wrote in the cover letter, it
is likely I'll propose similar C++ and Fortran directories in the
future.

Is the patch OK for trunk?

Thanks,

Martin


2016-03-03  Martin Jambor  <mjambor@suse.cz>

	* testsuite/lib/libgomp.exp
	(check_effective_target_hsa_offloading_selected_nocache): New.
	(check_effective_target_hsa_offloading_selected): Likewise.
	* testsuite/libgomp.hsa.c/c.exp: Likewise.
	* testsuite/libgomp.hsa.c/alloca-1.c: Likewise.
	* testsuite/libgomp.hsa.c/bitfield-1.c: Likewise.
	* testsuite/libgomp.hsa.c/builtins-1.c: Likewise.
	* testsuite/libgomp.hsa.c/complex-1.c: Likewise.
	* testsuite/libgomp.hsa.c/formal-actual-args-1.c: Likewise.
	* testsuite/libgomp.hsa.c/function-call-1.c: Likewise.
	* testsuite/libgomp.hsa.c/get-level-1.c: Likewise.
	* testsuite/libgomp.hsa.c/gridify-1.c: Likewise.
	* testsuite/libgomp.hsa.c/gridify-2.c: Likewise.
	* testsuite/libgomp.hsa.c/gridify-3.c: Likewise.
	* testsuite/libgomp.hsa.c/gridify-4.c: Likewise.
	* testsuite/libgomp.hsa.c/memory-operations-1.c: Likewise.
	* testsuite/libgomp.hsa.c/pr69568.c: Likewise.
	* testsuite/libgomp.hsa.c/rotate-1.c: Likewise.
	* testsuite/libgomp.hsa.c/switch-1.c: Likewise.
	* testsuite/libgomp.hsa.c/switch-branch-1.c: Likewise.
---
 libgomp/testsuite/lib/libgomp.exp                  |  53 +++++++
 libgomp/testsuite/libgomp.hsa.c/alloca-1.c         |  25 ++++
 libgomp/testsuite/libgomp.hsa.c/bitfield-1.c       | 160 +++++++++++++++++++++
 libgomp/testsuite/libgomp.hsa.c/builtins-1.c       |  97 +++++++++++++
 libgomp/testsuite/libgomp.hsa.c/c.exp              |  42 ++++++
 libgomp/testsuite/libgomp.hsa.c/complex-1.c        |  65 +++++++++
 .../testsuite/libgomp.hsa.c/formal-actual-args-1.c |  83 +++++++++++
 libgomp/testsuite/libgomp.hsa.c/function-call-1.c  |  50 +++++++
 libgomp/testsuite/libgomp.hsa.c/get-level-1.c      |  26 ++++
 libgomp/testsuite/libgomp.hsa.c/gridify-1.c        |  26 ++++
 libgomp/testsuite/libgomp.hsa.c/gridify-2.c        |  26 ++++
 libgomp/testsuite/libgomp.hsa.c/gridify-3.c        |  39 +++++
 libgomp/testsuite/libgomp.hsa.c/gridify-4.c        |  45 ++++++
 .../testsuite/libgomp.hsa.c/memory-operations-1.c  |  92 ++++++++++++
 libgomp/testsuite/libgomp.hsa.c/pr69568.c          |  41 ++++++
 libgomp/testsuite/libgomp.hsa.c/rotate-1.c         |  39 +++++
 libgomp/testsuite/libgomp.hsa.c/switch-1.c         | 145 +++++++++++++++++++
 libgomp/testsuite/libgomp.hsa.c/switch-branch-1.c  | 116 +++++++++++++++
 18 files changed, 1170 insertions(+)
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/alloca-1.c
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/bitfield-1.c
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/builtins-1.c
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/c.exp
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/complex-1.c
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/formal-actual-args-1.c
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/function-call-1.c
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/get-level-1.c
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/gridify-1.c
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/gridify-2.c
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/gridify-3.c
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/gridify-4.c
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/memory-operations-1.c
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/pr69568.c
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/rotate-1.c
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/switch-1.c
 create mode 100644 libgomp/testsuite/libgomp.hsa.c/switch-branch-1.c

diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index bbc2c26..0d5b6d4 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -395,3 +395,56 @@ proc check_effective_target_openacc_host_selected { } {
     }
     return 0;
 }
+
+# Return 1 if the selected OMP device is actually a HSA device
+
+proc check_effective_target_hsa_offloading_selected_nocache {} {
+    global tool
+
+    set src {
+	int main () {
+	    int v = 1;
+	    #pragma omp target map(from:v)
+	    v = 0;
+	    return v;
+	}
+    }
+    
+    set result [eval [list check_compile hsa_offloading_src executable $src] ""]
+    set lines [lindex $result 0]
+    set output [lindex $result 1]
+
+    set ok 0
+    if { [string match "" $lines] } {
+	# No error messages, let us switch on HSA debugging output and run it
+	set prev_HSA_DEBUG [getenv HSA_DEBUG]
+	setenv HSA_DEBUG "1"
+	set result [remote_load target "./$output" "2>&1" ""]
+	if { [string match "" $prev_HSA_DEBUG] } {
+	    unsetenv HSA_DEBUG
+	} else {
+	    setenv HSA_DEBUG $prev_HSA_DEBUG
+	}
+	set status [lindex $result 0]
+	if { $status != "pass" } {
+	    verbose "HSA availability test failed"
+	    return 0
+	}
+	set output [lindex $result 1]
+	if { [string match "*HSA debug: Going to dispatch kernel*" $output] } {
+	    verbose "HSA availability detected"
+	    set ok 1
+	}
+    }
+    remote_file build delete $output
+    return $ok
+}
+
+# Return 1 if the selected OMP device is actually a HSA device and
+# cache the result
+
+proc check_effective_target_hsa_offloading_selected {} {
+    return [check_cached_effective_target hsa_offloading_selected {
+	check_effective_target_hsa_offloading_selected_nocache
+    }]
+}
diff --git a/libgomp/testsuite/libgomp.hsa.c/alloca-1.c b/libgomp/testsuite/libgomp.hsa.c/alloca-1.c
new file mode 100644
index 0000000..48dca94
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/alloca-1.c
@@ -0,0 +1,25 @@
+#define size 10
+int i, j, k;
+
+int
+main ()
+{
+  char *s = __builtin_malloc (size + 1);
+
+#pragma omp target teams
+  {
+#pragma omp distribute parallel for default(none) private(i) shared(s)
+    for (i = 0; i < size; ++i)
+      {
+	char *buffer = __builtin_alloca (10);
+	buffer[5] = 97 + i;
+	s[i] = buffer[5];
+      }
+  }
+
+  for (i = 0; i < size; ++i)
+    if (s[i] != 97 + i)
+      __builtin_abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.hsa.c/bitfield-1.c b/libgomp/testsuite/libgomp.hsa.c/bitfield-1.c
new file mode 100644
index 0000000..4dbf348
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/bitfield-1.c
@@ -0,0 +1,160 @@
+#include <assert.h>
+
+#define ASSIGN_SX(N)                                                           \
+  s##N.a1 = 1;                                                                 \
+  s##N.a2 = 2;                                                                 \
+  s##N.a3 = 3;                                                                 \
+  s##N.a4 = 4;                                                                 \
+  s##N.a5 = 5;                                                                 \
+  s##N.a6 = 6;                                                                 \
+  s##N.a7 = 7;                                                                 \
+  s##N.a8 = 8;                                                                 \
+  s##N.a9 = 9;                                                                 \
+  s##N.a10 = 10;
+
+#define ASSERT_SX(N)                                                           \
+  assert (s##N.a1 == 1); \
+  assert (s##N.a2 == 2); \
+  assert (s##N.a3 == 3); \
+  assert (s##N.a4 == 4); \
+  assert (s##N.a5 == 5); \
+  assert (s##N.a6 == 6); \
+  assert (s##N.a7 == 7); \
+  assert (s##N.a8 == 8); \
+  assert (s##N.a9 == 9); \
+  assert (s##N.a10 == 10);
+
+struct S1
+{
+  unsigned a : 10;
+  unsigned b : 20;
+};
+
+struct S2
+{
+  unsigned a1 : 10;
+  unsigned a2 : 10;
+  unsigned a3 : 10;
+  unsigned a4 : 10;
+  unsigned a5 : 10;
+  unsigned a6 : 10;
+  unsigned a7 : 10;
+  unsigned a8 : 10;
+  unsigned a9 : 10;
+  unsigned a10 : 10;
+};
+
+struct S3
+{
+  unsigned a1 : 10;
+  unsigned a2 : 9;
+  unsigned a3 : 8;
+  unsigned a4 : 7;
+  unsigned a5 : 6;
+  unsigned a6 : 5;
+  unsigned a7 : 6;
+  unsigned a8 : 7;
+  unsigned a9 : 8;
+  unsigned a10 : 9;
+};
+
+struct S4
+{
+  unsigned a1 : 10;
+  int a2 : 9;
+  unsigned a3 : 8;
+  int a4 : 7;
+  unsigned a5 : 6;
+  int a6 : 5;
+  unsigned a7 : 6;
+  int a8 : 7;
+  unsigned a9 : 8;
+  int a10 : 9;
+};
+
+struct S5
+{
+  unsigned a1 : 31;
+  int a2 : 9;
+  unsigned a3 : 17;
+  int a4 : 7;
+  unsigned a5 : 6;
+  int a6 : 5;
+  unsigned long a7 : 55;
+  int a8 : 7;
+  unsigned a9 : 8;
+  int a10 : 9;
+};
+
+int
+main ()
+{
+  struct S1 s1;
+
+#pragma omp target map(to: s1)
+  {
+    s1.a = 2;
+    s1.b = 3;
+  }
+
+  assert (s1.a == 2);
+  assert (s1.b == 3);
+
+  struct S2 s2;
+
+#pragma omp target map(to: s2)
+  {
+    ASSIGN_SX (2)
+  }
+
+  ASSERT_SX (2)
+
+  struct S3 s3;
+
+#pragma omp target map(to: s3)
+  {
+    ASSIGN_SX (3)
+  }
+
+  ASSERT_SX (3)
+
+  struct S4 s4;
+
+#pragma omp target map(to: s4)
+  {
+    ASSIGN_SX (4)
+  }
+
+  ASSERT_SX (4)
+
+  struct S4 s5;
+
+  s5.a1 = 0;
+  s5.a2 = 1;
+  s5.a3 = 2;
+  s5.a4 = 3;
+  s5.a5 = 4;
+  s5.a6 = 5;
+  s5.a7 = 6;
+  s5.a8 = 7;
+  s5.a9 = 8;
+  s5.a10 = 9;
+
+#pragma omp target map(to: s5)
+  {
+    s5.a1++;
+    s5.a2++;
+    s5.a3++;
+    s5.a4++;
+    s5.a5++;
+    s5.a6++;
+    s5.a7++;
+    s5.a8++;
+    s5.a9++;
+    s5.a10++;
+  }
+
+  ASSERT_SX (5)
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.hsa.c/builtins-1.c b/libgomp/testsuite/libgomp.hsa.c/builtins-1.c
new file mode 100644
index 0000000..e603c21
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/builtins-1.c
@@ -0,0 +1,97 @@
+/* { dg-additional-options "-ffast-math" } */
+
+#include <assert.h>
+#include <math.h>
+
+#define N 10
+#define N2 14
+
+#define c1 1.2345f
+#define c2 1.2345
+
+#define DELTA 0.001
+
+#define TEST_BIT_BUILTINS(T, S, S2)                                            \
+  {                                                                            \
+    T arguments[N2]                                                            \
+      = {0##S,		1##S,	  2##S,	  3##S,                    \
+	 111##S,	333##S,	444##S,	0x80000000##S,           \
+	 0x0000ffff##S, 0xf0000000##S, 0xff000000##S, 0xffffffff##S};          \
+    int clrsb[N2] = {};                                                        \
+    int clz[N2] = {};                                                          \
+    int ctz[N2] = {};                                                          \
+    int ffs[N2] = {};                                                          \
+    int parity[N2] = {};                                                       \
+    int popcount[N2] = {};                                                     \
+                                                                               \
+    _Pragma ("omp target map(to:clz[:N2], ctz[:N2], ffs[:N2], parity[:N2], popcount[:N2])")                                                 \
+    {                                                                          \
+      for (unsigned i = 0; i < N2; i++)                                        \
+	{                                                                      \
+	  clrsb[i] = __builtin_clrsb##S2 (arguments[i]);                       \
+	  clz[i] = __builtin_clz##S2 (arguments[i]);                           \
+	  ctz[i] = __builtin_ctz##S2 (arguments[i]);                           \
+	  ffs[i] = __builtin_ffs##S2 (arguments[i]);                           \
+	  parity[i] = __builtin_parity##S2 (arguments[i]);                     \
+	  popcount[i] = __builtin_popcount##S2 (arguments[i]);                 \
+	}                                                                      \
+    }                                                                          \
+                                                                               \
+    for (unsigned i = 0; i < N2; i++)                                          \
+      {                                                                        \
+	assert (clrsb[i] == __builtin_clrsb##S2 (arguments[i]));               \
+	if (arguments[0] != 0)                                                 \
+	  {                                                                    \
+	    assert (clz[i] == __builtin_clz##S2 (arguments[i]));               \
+	    assert (ctz[i] == __builtin_ctz##S2 (arguments[i]));               \
+	  }                                                                    \
+	assert (ffs[i] == __builtin_ffs##S2 (arguments[i]));                   \
+	assert (parity[i] == __builtin_parity##S2 (arguments[i]));             \
+	assert (popcount[i] == __builtin_popcount##S2 (arguments[i]));         \
+      }                                                                        \
+  }
+
+#define ASSERT(v1, v2) assert (fabs (v1 - v2) < DELTA)
+
+int
+main ()
+{
+  float f[N] = {};
+  float d[N] = {};
+
+/* 1) test direct mapping to HSA insns.  */
+
+#pragma omp target map(to: f[ : N], d[ : N])
+  {
+    f[0] = sinf (c1);
+    f[1] = cosf (c1);
+    f[2] = exp2f (c1);
+    f[3] = log2f (c1);
+    f[4] = truncf (c1);
+    f[5] = sqrtf (c1);
+
+    d[0] = trunc (c2);
+    d[1] = sqrt (c2);
+  }
+
+  ASSERT (f[0], sinf (c1));
+  ASSERT (f[1], cosf (c1));
+  ASSERT (f[2], exp2f (c1));
+  ASSERT (f[3], log2f (c1));
+  ASSERT (f[4], truncf (c1));
+  ASSERT (f[5], sqrtf (c1));
+
+  ASSERT (d[0], trunc (c2));
+  ASSERT (d[1], sqrt (c2));
+
+  /* 2) test bit builtins for unsigned int.  */
+  TEST_BIT_BUILTINS (int, , );
+
+  /* 3) test bit builtins for unsigned long int.  */
+  TEST_BIT_BUILTINS (long, l, l);
+
+  /* 4) test bit builtins for unsigned long long int.  */
+  TEST_BIT_BUILTINS (long long, ll, ll);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.hsa.c/c.exp b/libgomp/testsuite/libgomp.hsa.c/c.exp
new file mode 100644
index 0000000..4614192
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/c.exp
@@ -0,0 +1,42 @@
+if [info exists lang_library_path] then {
+    unset lang_library_path
+    unset lang_link_flags
+}
+if [info exists lang_test_file] then {
+    unset lang_test_file
+}
+if [info exists lang_include_flags] then {
+    unset lang_include_flags
+}
+
+load_lib libgomp-dg.exp
+load_gcc_lib gcc-dg.exp
+
+# Initialize dg.
+dg-init
+
+# Turn on OpenMP.
+lappend ALWAYS_CFLAGS "additional_flags=-fopenmp"
+
+set ld_library_path $always_ld_library_path
+append ld_library_path [gcc-set-multilib-library-path $GCC_UNDER_TEST]
+set_ld_library_path_env_vars
+
+global DEFAULT_CFLAGS
+if [info exists DEFAULT_CFLAGS] then {
+    set CFLAGS_list [list "-O0" $DEFAULT_CFLAGS]
+} else {
+    set CFLAGS_list [list "-O0" "-O2"]
+}
+
+if [check_effective_target_hsa_offloading_selected] {
+    foreach USE_CFLAGS $CFLAGS_list {
+	# Gather a list of all tests.
+	set tests [lsort [find $srcdir/$subdir *.c]]
+	# Main loop.
+	dg-runtest $tests "" [concat $USE_CFLAGS "-Whsa"]
+    }
+}
+
+# All done.
+dg-finish
diff --git a/libgomp/testsuite/libgomp.hsa.c/complex-1.c b/libgomp/testsuite/libgomp.hsa.c/complex-1.c
new file mode 100644
index 0000000..438c64a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/complex-1.c
@@ -0,0 +1,65 @@
+#include <assert.h>
+#include <complex.h>
+#include <math.h>
+
+#define uchar unsigned char
+#define C 123
+
+#define TEST(type)                                                             \
+  type foo_##type (void)                                                       \
+  {                                                                            \
+    _Complex type a = C + 45I;                                                 \
+    return __real__ a;                                                         \
+  }
+
+#pragma omp declare target
+TEST (char)
+TEST (uchar)
+TEST (short)
+TEST (int)
+
+float
+bar (float a, float b)
+{
+  _Complex float c = a + b * I;
+
+  c += 11.f + 12.f * I;
+
+  _Complex float d = 2.f + 4.44f * I;
+
+  return __real__(crealf (c + d) + cimag (d) * I);
+}
+
+#pragma omp end declare target
+
+int
+main (void)
+{
+  int v = 0;
+  float v2 = 0.0f;
+
+#pragma omp target map(to: v)
+  v = foo_char ();
+
+  assert (v == C);
+
+#pragma omp target map(to: v)
+  v = foo_uchar ();
+
+  assert (v == C);
+
+#pragma omp target map(to: v)
+  v = foo_short ();
+
+  assert (v == C);
+
+#pragma omp target map(to: v)
+  v = foo_int ();
+
+  assert (v == C);
+
+#pragma omp target map(to: v2)
+  v2 = bar (1.12f, 4.44f);
+
+  assert (fabs (v2 - 14.12) < 0.0001f);
+}
diff --git a/libgomp/testsuite/libgomp.hsa.c/formal-actual-args-1.c b/libgomp/testsuite/libgomp.hsa.c/formal-actual-args-1.c
new file mode 100644
index 0000000..058a036
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/formal-actual-args-1.c
@@ -0,0 +1,83 @@
+#include <assert.h>
+
+struct Cube
+{
+  int x;
+  int y;
+  int z;
+};
+
+#pragma omp declare target
+int
+foo (short a)
+{
+  switch (a)
+    {
+    case 1:
+      return 11;
+      break;
+    case 33:
+      return 333;
+      break;
+    case 55:
+      return 55;
+      break;
+    default:
+      return -1;
+    }
+}
+
+int
+bar (int a)
+{
+  int *ptr = &a;
+
+  *ptr = 100;
+  return a + *ptr;
+}
+
+struct Cube
+baz (struct Cube c)
+{
+  c.x = 11;
+  return c;
+}
+
+#pragma omp end declare target
+
+#define s 100
+
+int
+main (int argc)
+{
+  /* Test 1: argument types: char to short.  */
+
+  int array[s];
+#pragma omp target map(tofrom : array[ : s])
+  {
+    for (char i = 0; i < s; i++)
+      array[i] = foo (i);
+  }
+
+  for (int i = 0; i < s; i++)
+    assert (array[i] == foo (i));
+
+  /* Test 2: argument address is taken.  */
+  int v = 2;
+
+#pragma omp target map(tofrom : v)
+  v = bar (v);
+
+  assert (v == 200);
+
+  /* Test 3: passing a structure as a function argument.  */
+  struct Cube r;
+  struct Cube c = {.x = 1, .y = 2, .z = 3};
+
+#pragma omp target map(to : r) map(from : c)
+  r = baz (c);
+
+  assert (r.x == 11);
+  assert (r.y == c.y);
+  assert (r.z == c.z);
+}
diff --git a/libgomp/testsuite/libgomp.hsa.c/function-call-1.c b/libgomp/testsuite/libgomp.hsa.c/function-call-1.c
new file mode 100644
index 0000000..7f15dff
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/function-call-1.c
@@ -0,0 +1,50 @@
+#define size 8
+
+#pragma omp declare target
+int
+identity (int x)
+{
+  return x;
+}
+
+int
+expx (int x, int n)
+{
+  for (int i = 0; i < n - 1; i++)
+    x *= x;
+
+  return x;
+}
+
+float
+init (int x, int y)
+{
+  int x1 = identity (identity (identity (identity (x))));
+  int y1 = identity (identity (identity (identity (y))));
+
+  int x2 = expx (x1, 2);
+  int y2 = expx (y1, 2);
+
+  return (x2 + y2);
+}
+#pragma omp end declare target
+
+int
+main ()
+{
+  int i, j;
+  int a[size][size];
+
+#pragma omp target teams map(to:a[:size][:size])
+#pragma omp distribute parallel for default(none) private(i, j) shared(a)
+  for (i = 0; i < size; ++i)
+    for (j = 0; j < size; ++j)
+      a[i][j] = init (i, j);
+
+  for (i = 0; i < size; ++i)
+    for (j = 0; j < size; ++j)
+      if (i * i + j * j != a[i][j])
+       __builtin_abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.hsa.c/get-level-1.c b/libgomp/testsuite/libgomp.hsa.c/get-level-1.c
new file mode 100644
index 0000000..81c9df0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/get-level-1.c
@@ -0,0 +1,26 @@
+#include <omp.h>
+
+int
+main ()
+{
+  int i;
+  int level = -1;
+
+#pragma omp target map(tofrom : level)
+  {
+    level = omp_get_level ();
+  }
+
+  if (level != 0)
+    __builtin_abort ();
+
+#pragma omp target teams map(tofrom : level)
+#pragma omp distribute parallel for default(none) private(i) shared(level)
+  for (i = 0; i < 1; ++i)
+    level += omp_get_level ();
+
+  if (level != 1)
+    __builtin_abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.hsa.c/gridify-1.c b/libgomp/testsuite/libgomp.hsa.c/gridify-1.c
new file mode 100644
index 0000000..b670b9b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/gridify-1.c
@@ -0,0 +1,26 @@
+void __attribute__((noinline, noclone))
+foo (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]++;
+}
+
+int main (int argc, char **argv)
+{
+  int n = 32;
+  int *a = __builtin_malloc (sizeof (int) * n);
+  int i;
+
+  __builtin_memset (a, 0, sizeof (int) * n);
+  foo (n, a, 32);
+  for (i = 0; i < n; i ++)
+    {
+      if (a[i] != 1)
+	__builtin_abort ();
+    }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.hsa.c/gridify-2.c b/libgomp/testsuite/libgomp.hsa.c/gridify-2.c
new file mode 100644
index 0000000..3692eb0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/gridify-2.c
@@ -0,0 +1,26 @@
+void __attribute__((noinline, noclone))
+foo (int j, int n, int *a)
+{
+  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++)
+      a[i] = i;
+}
+
+int main (int argc, char **argv)
+{
+  int n = 32;
+  int *a = __builtin_malloc (sizeof (int) * n);
+  int i, j = 4;
+
+  __builtin_memset (a, 0, sizeof (int) * n);
+  foo (j, n, a);
+  for (i = j + 1; i < n; i ++)
+    {
+      if (a[i] != i)
+	__builtin_abort ();
+    }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.hsa.c/gridify-3.c b/libgomp/testsuite/libgomp.hsa.c/gridify-3.c
new file mode 100644
index 0000000..f881d81
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/gridify-3.c
@@ -0,0 +1,39 @@
+#define THE_LOOP \
+  for (i = j + 1; i < n; i += 3) \
+    a[i] = i
+
+void __attribute__((noinline, noclone))
+foo (int j, int n, int *a)
+{
+  int i;
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for shared(a) firstprivate(n) private(i) firstprivate(j)
+  THE_LOOP;
+}
+
+void __attribute__((noinline, noclone))
+bar (int j, int n, int *a)
+{
+  int i;
+  THE_LOOP;
+}
+
+int main (int argc, char **argv)
+{
+  int n = 32;
+  int *a = __builtin_malloc (sizeof (int) * n);
+  int *ref = __builtin_malloc (sizeof (int) * n);
+  int i, j = 4;
+
+  __builtin_memset (a, 0, sizeof (int) * n);
+  __builtin_memset (ref, 0, sizeof (int) * n);
+  bar (j, n, ref);
+  foo (j, n, a);
+  for (i = 0; i < n; i ++)
+    {
+      if (a[i] != ref[i])
+	__builtin_abort ();
+    }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.hsa.c/gridify-4.c b/libgomp/testsuite/libgomp.hsa.c/gridify-4.c
new file mode 100644
index 0000000..c3fbdbf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/gridify-4.c
@@ -0,0 +1,45 @@
+#define THE_LOOP \
+  for (i = j + 1; i < n; i += 3) \
+    a[i] = i
+
+void __attribute__((noinline, noclone))
+foo (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)
+      THE_LOOP;
+    }
+  }
+}
+
+void __attribute__((noinline, noclone))
+bar (int j, int n, int *a)
+{
+  int i;
+  THE_LOOP;
+}
+
+int main (int argc, char **argv)
+{
+  int n = 32;
+  int *a = __builtin_malloc (sizeof (int) * n);
+  int *ref = __builtin_malloc (sizeof (int) * n);
+  int i, j = 4;
+
+  __builtin_memset (a, 0, sizeof (int) * n);
+  __builtin_memset (ref, 0, sizeof (int) * n);
+  bar (j, n, ref);
+  foo (j, n, a);
+  for (i = 0; i < n; i ++)
+    {
+      if (a[i] != ref[i])
+	__builtin_abort ();
+    }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.hsa.c/memory-operations-1.c b/libgomp/testsuite/libgomp.hsa.c/memory-operations-1.c
new file mode 100644
index 0000000..a17be93
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/memory-operations-1.c
@@ -0,0 +1,92 @@
+#include <assert.h>
+
+#define C 55
+
+int i, j, k;
+
+static void
+test_bzero (unsigned size)
+{
+  unsigned bsize = size * sizeof (int);
+  int *x = __builtin_malloc (bsize);
+  __builtin_memset (x, C, bsize);
+
+#pragma omp target map(tofrom: x[:size]) map(from: bsize)
+  {
+    __builtin_bzero (x, bsize);
+  }
+
+  char *buffer = (char *) x;
+  for (unsigned i = 0; i < bsize; ++i)
+    assert (buffer[i] == 0);
+}
+
+static void
+test_memcpy (unsigned size)
+{
+  unsigned bsize = size * sizeof (int);
+  int *x = __builtin_malloc (bsize);
+  __builtin_memset (x, C, bsize);
+  int *y = __builtin_malloc (bsize);
+
+#pragma omp target map(tofrom: x[:size], y[:size]) map(from: bsize)
+  {
+    __builtin_memcpy (y, x, bsize);
+  }
+
+  char *buffer = (char *) y;
+  for (unsigned i = 0; i < bsize; ++i)
+    assert (buffer[i] == C);
+}
+
+static void
+test_mempcpy (unsigned size)
+{
+  unsigned bsize = size * sizeof (int);
+  int *x = __builtin_malloc (bsize);
+  __builtin_memset (x, C, bsize);
+  int *y = __builtin_malloc (bsize);
+  int *ptr = 0;
+
+#pragma omp target map(tofrom :x[:size], y[:size], ptr) map(from: bsize)
+  {
+    ptr = __builtin_mempcpy (y, x, bsize);
+  }
+
+  char *buffer = (char *) y;
+  for (unsigned i = 0; i < bsize; ++i)
+    assert (buffer[i] == C);
+
+  assert (ptr == y + size);
+}
+
+static void
+test_memset (unsigned size)
+{
+  unsigned bsize = size * sizeof (int);
+  int *x = __builtin_malloc (bsize);
+  __builtin_bzero (x, bsize);
+
+#pragma omp target map(tofrom : x[:size]) map(from: bsize)
+  {
+    __builtin_memset (x, C, bsize);
+  }
+
+  char *buffer = (char *) x;
+  for (unsigned i = 0; i < bsize; ++i)
+    assert (buffer[i] == C);
+}
+
+int
+main (void)
+{
+  unsigned tests[] = {1, 2, 3, 4, 5, 8, 15, 17, 23, 33, 0};
+
+  for (unsigned i = 0; tests[i]; i++)
+    {
+      test_bzero (tests[i]);
+      test_memset (tests[i]);
+      test_memcpy (tests[i]);
+      test_mempcpy (tests[i]);
+    }
+}
diff --git a/libgomp/testsuite/libgomp.hsa.c/pr69568.c b/libgomp/testsuite/libgomp.hsa.c/pr69568.c
new file mode 100644
index 0000000..6262eee
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/pr69568.c
@@ -0,0 +1,41 @@
+/* PR hsa/69568 */
+
+typedef float float2 __attribute__ ((vector_size (8)));
+float2 *output;
+
+void __attribute__((noinline, noclone))
+foo (int n, float2 *a, int workgroup_size)
+{
+  int i;
+#pragma omp target map(from:a[:n]) firstprivate(n, workgroup_size)
+#pragma omp teams thread_limit(workgroup_size)
+#pragma omp distribute parallel for shared(a) firstprivate(n) private(i)
+    for (i = 0; i < n; i++)
+      { float2 v;
+	v[0] = i;
+	v[1] = 1+i;
+	a[i] = v;
+      }
+}
+
+int main (int argc, char **argv)
+{
+  int n = 32;
+  float2 *a = __builtin_malloc (sizeof (float2) * n);
+  int i;
+
+  __builtin_memset (a, 0, sizeof (float2) * n);
+  foo (n, a, 32);
+  for (i = 0; i < n; i++)
+    {
+      float2 v = a[i];
+      if (__builtin_abs (v[0] - i) > 0.1
+	  || __builtin_abs (v[1] - i - 1) > 0.1)
+	{
+	  __builtin_abort ();
+	  return 1;
+	}
+    }
+  return 0;
+}
+
diff --git a/libgomp/testsuite/libgomp.hsa.c/rotate-1.c b/libgomp/testsuite/libgomp.hsa.c/rotate-1.c
new file mode 100644
index 0000000..494388b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/rotate-1.c
@@ -0,0 +1,39 @@
+#include <assert.h>
+#include <limits.h>
+
+#define T unsigned int
+#define BITSIZE CHAR_BIT * sizeof (T)
+
+#define C1 123u
+
+#pragma omp declare target
+T
+rotate (T value, T shift)
+{
+  T r = (value << shift) | (value >> (BITSIZE - shift));
+  return (r >> shift) | (r << (BITSIZE - shift));
+}
+#pragma omp end declare target
+
+int
+main (int argc)
+{
+  T v1, v2, v3, v4, v5;
+
+#pragma omp target map(to: v1, v2, v3, v4, v5)
+  {
+    v1 = rotate (C1, 10);
+    v2 = rotate (C1, 2);
+    v3 = rotate (C1, 5);
+    v4 = rotate (C1, 16);
+    v5 = rotate (C1, 32);
+  }
+
+  assert (v1 == C1);
+  assert (v2 == C1);
+  assert (v3 == C1);
+  assert (v4 == C1);
+  assert (v5 == C1);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.hsa.c/switch-1.c b/libgomp/testsuite/libgomp.hsa.c/switch-1.c
new file mode 100644
index 0000000..a180cf6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/switch-1.c
@@ -0,0 +1,145 @@
+#include <assert.h>
+
+#define s 100
+
+#pragma omp declare target
+int
+switch1 (int a)
+{
+  switch (a)
+    {
+    case 1:
+      return 11;
+    case 33:
+      return 333;
+    case 55:
+      return 55;
+    default:
+      return -1;
+    }
+}
+
+int
+switch2 (int a)
+{
+  switch (a)
+    {
+    case 1 ... 11:
+      return 11;
+      break;
+    case 33:
+      return 333;
+      break;
+    case 55:
+      return 55;
+      break;
+    default:
+      return -1;
+    }
+}
+
+int
+switch3 (int a)
+{
+  switch (a)
+    {
+    case 1 ... 11:
+      return 11;
+    case 12 ... 22:
+      return 22;
+    case 23 ... 33:
+      return 33;
+    case 34 ... 44:
+      return 44;
+    default:
+      return 44;
+    }
+}
+
+int
+switch4 (int a, int b)
+{
+  switch (a)
+    {
+    case 1 ... 11:
+      return a;
+    case 12 ... 22:
+      return b;
+    case 23 ... 33:
+      return a;
+    case 34 ... 44:
+      return b;
+    default:
+      return 12345;
+    }
+}
+
+int
+switch5 (int a, int b)
+{
+  switch (a)
+    {
+    case 1 ... 2:
+      return 1;
+    case 3 ... 4:
+      return 2;
+    case 5 ... 6:
+      return 3;
+    case 7 ... 11:
+      return 4;
+    }
+
+  return -1;
+}
+#pragma omp end declare target
+
+int
+main (int argc)
+{
+  int array[s];
+
+#pragma omp target map(tofrom : array[:s])
+  {
+    for (int i = 0; i < s; i++)
+      array[i] = switch1 (i);
+  }
+
+  for (int i = 0; i < s; i++)
+    assert (array[i] == switch1 (i));
+
+#pragma omp target map(tofrom : array[:s])
+  {
+    for (int i = 0; i < s; i++)
+      array[i] = switch2 (i);
+  }
+
+  for (int i = 0; i < s; i++)
+    assert (array[i] == switch2 (i));
+
+#pragma omp target map(tofrom : array[:s])
+  {
+    for (int i = 0; i < s; i++)
+      array[i] = switch3 (i);
+  }
+
+  for (int i = 0; i < s; i++)
+    assert (array[i] == switch3 (i));
+
+#pragma omp target map(tofrom : array[:s])
+  {
+    for (int i = 0; i < s; i++)
+      array[i] = switch4 (i, i + 1);
+  }
+
+  for (int i = 0; i < s; i++)
+    assert (array[i] == switch4 (i, i + 1));
+
+#pragma omp target map(tofrom : array[:s])
+  {
+    for (int i = 0; i < s; i++)
+      array[i] = switch5 (i, i + 1);
+  }
+
+  for (int i = 0; i < s; i++)
+    assert (array[i] == switch5 (i, i + 1));
+}
diff --git a/libgomp/testsuite/libgomp.hsa.c/switch-branch-1.c b/libgomp/testsuite/libgomp.hsa.c/switch-branch-1.c
new file mode 100644
index 0000000..9af1d6d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/switch-branch-1.c
@@ -0,0 +1,116 @@
+#include <assert.h>
+
+#define s 100
+
+#pragma omp declare target
+int
+switch1 (unsigned a)
+{
+  switch (a)
+    {
+    case 1 ... 11:
+      return 11;
+    case 12 ... 13:
+      return 22;
+    default:
+      return 44;
+    }
+}
+
+int
+switch2 (unsigned a)
+{
+  switch (a)
+    {
+    case 1 ... 5:
+      return 1;
+    case 9 ... 11:
+      return a + 3;
+    case 12 ... 13:
+      return a + 3;
+    default:
+      return 44;
+    }
+}
+
+#define OFFSET 12
+
+int
+switch3 (unsigned a)
+{
+  switch (a)
+    {
+    case (OFFSET + 0):
+      return 1;
+    case (OFFSET + 1)...(OFFSET + 11):
+      return 11;
+    case (OFFSET + 12)...(OFFSET + 13):
+      return (OFFSET + 22);
+    default:
+      return (OFFSET + 44);
+    }
+}
+
+int
+switch4 (unsigned a)
+{
+  switch (a)
+    {
+    case -2:
+      return 1;
+    case -1:
+      return a + 3;
+    case 3:
+      return a + 3;
+    default:
+      return 44;
+    }
+}
+#pragma omp end declare target
+
+#define low -33
+#define high 55
+
+int
+main (int argc)
+{
+  int array[s];
+
+#pragma omp target map(tofrom : array[:s])
+  {
+    for (int i = low; i < high; i++)
+      array[i - low] = switch1 (i);
+  }
+
+  for (int i = low; i < high; i++)
+    assert (array[i - low] == switch1 (i));
+
+#pragma omp target map(tofrom : array[:s])
+  {
+    for (int i = low; i < high; i++)
+      array[i - low] = switch2 (i);
+  }
+
+  for (int i = low; i < high; i++)
+    assert (array[i - low] == switch2 (i));
+
+#pragma omp target map(tofrom : array[:s])
+  {
+    for (int i = low; i < high; i++)
+      array[i - low] = switch3 (i);
+  }
+
+  for (int i = low; i < high; i++)
+    assert (array[i - low] == switch3 (i));
+
+#pragma omp target map(tofrom : array[:s])
+  {
+    for (int i = low; i < high; i++)
+      array[i - low] = switch4 (i);
+  }
+
+  for (int i = low; i < high; i++)
+    assert (array[i - low] == switch4 (i));
+
+  return 0;
+}
-- 
2.7.1

[hsa testsuite 0/5] Re-post of all pending patches adjusting testsuite for HSA

[Martin Jambor]

Hi,

in order to consolidate things, I have decided to re-post all "hsa
testsuite" patches under this thread.  With the patches applied, we do
no not get any spurious failures because of hsa warnings or libgomp
testcases failing because they are run on the host fallback.
Moreover, the first patch adds a simple dump-scan compile-time
gridification tests and the last patch adds a special directory for
run-time C tests of hsa which are run only when HSA devices are
actually selected for offloading.  In the future, I'll likely propose
similar C++ and Fortran directories.

All patches were tested by running the whole testsuite on patched
trunk:

  - that was configured for all languages except go but not configured
    for HSA,

  - that was configured for all languages except go and also for HSA
    offloading, but an HSA device was not present on the machine, and

  - running the whole suite after configuring trunk for C, C++ and
    Fortran on a computer with an HSA APU,

and subsequently comparing generated .sum files with unpatched trunk.

Thanks for any feedback (and approvals ;-),

Martin

[hsa testsuite 2/5] Suppress hsa warnings in compiler gomp tests

[Martin Jambor]

Hi,

as Jakub requested, this patch deals with HSA "excess errors" in the
gomp compiler testsuite by passing -Wno-hsa to all of them.  After
discussing this in the thread about similar libgomp tests[1] (which
are however handled differently), Jakub expressed preference for
passing the option in default_extra_flags rather than flags so that
names of the tests do not change.

This however requires that the failing tests which use dg-options must
be adjusted.  There is 9 of them, most of them have just superfluous
-fopenmp in them which can be removed because that is the default and
the rest is handled by turning dg-options into dg-additional-options.

OK for trunk?

Thanks,

Martin

[1] https://gcc.gnu.org/ml/gcc-patches/2016-03/msg00381.html


2016-03-04  Martin Jambor  <mjambor@suse.cz>

	* c-c++-common/gomp/clauses-1.c: Remove dg-options.
	* c-c++-common/gomp/if-1.c: Likewise.
	* c-c++-common/gomp/pr61486-2.c: Likewise.
	* c-c++-common/gomp/target-teams-1.c: Moved dg-options except -fopenmp
	to dg-additional-options.
	* g++.dg/gomp/gomp.exp: Pass -Wno-hsa to all tests.
	* g++/gomp/target-teams-1.c: Likewise.
	* gcc.dg/gomp/gomp.exp: Likewise.
	* gcc.dg/gomp/pr68128-2.c: Moved dg-options except -fopenmp to
	dg-additional-options.
	* gfortran.dg/gomp/gomp.exp: Likewise.
	* gfortran.dg/gomp/target1.f90: Remove dg-options.
	* gfortran.dg/gomp/target2.f90: Moved dg-options except -fopenmp to
	dg-additional-options.
	* gfortran.dg/gomp/target3.f90: Remove dg-options.
---
 gcc/testsuite/c-c++-common/gomp/clauses-1.c      | 1 -
 gcc/testsuite/c-c++-common/gomp/if-1.c           | 1 -
 gcc/testsuite/c-c++-common/gomp/pr61486-2.c      | 1 -
 gcc/testsuite/c-c++-common/gomp/target-teams-1.c | 2 +-
 gcc/testsuite/g++.dg/gomp/gomp.exp               | 2 +-
 gcc/testsuite/g++.dg/gomp/target-teams-1.C       | 2 +-
 gcc/testsuite/gcc.dg/gomp/gomp.exp               | 2 +-
 gcc/testsuite/gcc.dg/gomp/pr68128-2.c            | 2 +-
 gcc/testsuite/gfortran.dg/gomp/gomp.exp          | 2 +-
 gcc/testsuite/gfortran.dg/gomp/target1.f90       | 1 -
 gcc/testsuite/gfortran.dg/gomp/target2.f90       | 2 +-
 gcc/testsuite/gfortran.dg/gomp/target3.f90       | 1 -
 12 files changed, 7 insertions(+), 12 deletions(-)

diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-1.c b/gcc/testsuite/c-c++-common/gomp/clauses-1.c
index 2d1c352..91aed39 100644
--- a/gcc/testsuite/c-c++-common/gomp/clauses-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/clauses-1.c
@@ -1,5 +1,4 @@
 /* { dg-do compile } */
-/* { dg-options "-fopenmp" } */
 /* { dg-additional-options "-std=c99" { target c } } */
 
 int t;
diff --git a/gcc/testsuite/c-c++-common/gomp/if-1.c b/gcc/testsuite/c-c++-common/gomp/if-1.c
index 4ba708c..3a9b538 100644
--- a/gcc/testsuite/c-c++-common/gomp/if-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/if-1.c
@@ -1,5 +1,4 @@
 /* { dg-do compile } */
-/* { dg-options "-fopenmp" } */
 
 void
 foo (int a, int b, int *p, int *q)
diff --git a/gcc/testsuite/c-c++-common/gomp/pr61486-2.c b/gcc/testsuite/c-c++-common/gomp/pr61486-2.c
index db97143..4a68023 100644
--- a/gcc/testsuite/c-c++-common/gomp/pr61486-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/pr61486-2.c
@@ -1,6 +1,5 @@
 /* PR middle-end/61486 */
 /* { dg-do compile } */
-/* { dg-options "-fopenmp" } */
 /* { dg-require-effective-target alloca } */
 
 #pragma omp declare target
diff --git a/gcc/testsuite/c-c++-common/gomp/target-teams-1.c b/gcc/testsuite/c-c++-common/gomp/target-teams-1.c
index 0a707c2..51b8d48 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-teams-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-teams-1.c
@@ -1,5 +1,5 @@
 /* { dg-do compile } */
-/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
 
 int v = 6;
 void bar (int);
diff --git a/gcc/testsuite/g++.dg/gomp/gomp.exp b/gcc/testsuite/g++.dg/gomp/gomp.exp
index 7365389..d26596c 100644
--- a/gcc/testsuite/g++.dg/gomp/gomp.exp
+++ b/gcc/testsuite/g++.dg/gomp/gomp.exp
@@ -29,7 +29,7 @@ dg-init
 # Main loop.
 g++-dg-runtest [lsort [concat \
 	[find $srcdir/$subdir *.C] \
-	[find $srcdir/c-c++-common/gomp *.c]]] "" "-fopenmp"
+	[find $srcdir/c-c++-common/gomp *.c]]] "" "-fopenmp -Wno-hsa"
 
 # All done.
 dg-finish
diff --git a/gcc/testsuite/g++.dg/gomp/target-teams-1.C b/gcc/testsuite/g++.dg/gomp/target-teams-1.C
index 0a97de0..f78a608 100644
--- a/gcc/testsuite/g++.dg/gomp/target-teams-1.C
+++ b/gcc/testsuite/g++.dg/gomp/target-teams-1.C
@@ -1,5 +1,5 @@
 // { dg-do compile }
-// { dg-options "-fopenmp -fdump-tree-gimple" }
+// { dg-additional-options "-fdump-tree-gimple" }
 
 int v = 6;
 void bar (int);
diff --git a/gcc/testsuite/gcc.dg/gomp/gomp.exp b/gcc/testsuite/gcc.dg/gomp/gomp.exp
index 78623fc..b6b5932 100644
--- a/gcc/testsuite/gcc.dg/gomp/gomp.exp
+++ b/gcc/testsuite/gcc.dg/gomp/gomp.exp
@@ -31,7 +31,7 @@ dg-init
 # Main loop.
 dg-runtest [lsort [concat \
 	[find $srcdir/$subdir *.c] \
-	[find $srcdir/c-c++-common/gomp *.c]]] "" "-fopenmp"
+	[find $srcdir/c-c++-common/gomp *.c]]] "" "-fopenmp -Wno-hsa"
 
 # All done.
 dg-finish
diff --git a/gcc/testsuite/gcc.dg/gomp/pr68128-2.c b/gcc/testsuite/gcc.dg/gomp/pr68128-2.c
index 58a07e9..9720add 100644
--- a/gcc/testsuite/gcc.dg/gomp/pr68128-2.c
+++ b/gcc/testsuite/gcc.dg/gomp/pr68128-2.c
@@ -1,6 +1,6 @@
 /* PR tree-optimization/68128 */
 /* { dg-do compile } */
-/* { dg-options "-O2 -fopenmp -fdump-tree-omplower" } */
+/* { dg-additional-options "-O2 -fdump-tree-omplower" } */
 
 extern int omp_get_thread_num (void);
 extern int omp_get_ancestor_thread_num (int);
diff --git a/gcc/testsuite/gfortran.dg/gomp/gomp.exp b/gcc/testsuite/gfortran.dg/gomp/gomp.exp
index 625361b..1cd2e36 100644
--- a/gcc/testsuite/gfortran.dg/gomp/gomp.exp
+++ b/gcc/testsuite/gfortran.dg/gomp/gomp.exp
@@ -30,7 +30,7 @@ dg-init
 
 # Main loop.
 gfortran-dg-runtest [lsort \
-       [find $srcdir/$subdir *.\[fF\]{,90,95,03,08} ] ] "" "-fopenmp"
+       [find $srcdir/$subdir *.\[fF\]{,90,95,03,08} ] ] "" "-fopenmp -Wno-hsa"
 
 # All done.
 dg-finish
diff --git a/gcc/testsuite/gfortran.dg/gomp/target1.f90 b/gcc/testsuite/gfortran.dg/gomp/target1.f90
index 14db497..1e77176 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target1.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target1.f90
@@ -1,5 +1,4 @@
 ! { dg-do compile }
-! { dg-options "-fopenmp" }
 
 module target1
   interface
diff --git a/gcc/testsuite/gfortran.dg/gomp/target2.f90 b/gcc/testsuite/gfortran.dg/gomp/target2.f90
index 7521331..dfe0ec3 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target2.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target2.f90
@@ -1,5 +1,5 @@
 ! { dg-do compile }
-! { dg-options "-fopenmp -ffree-line-length-160" }
+! { dg-additional-options "-ffree-line-length-160" }
 
 subroutine foo (n, s, t, u, v, w)
   integer :: n, i, s, t, u, v, w
diff --git a/gcc/testsuite/gfortran.dg/gomp/target3.f90 b/gcc/testsuite/gfortran.dg/gomp/target3.f90
index 7ba42a0..d968e54 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target3.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target3.f90
@@ -1,5 +1,4 @@
 ! { dg-do compile }
-! { dg-options "-fopenmp" }
 
 subroutine foo (r)
   integer :: i, r
-- 
2.7.1

Re: [hsa testsuite 1/5] Gridification tests

[Jakub Jelinek]

On Mon, Feb 29, 2016 at 04:55:44PM +0100, Martin Jambor wrote:
> 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?

Ok.
> 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.

	Jakub

Re: [hsa testsuite 3/5] Suppress hsa warnings in libgomp tests

[Jakub Jelinek]

On Mon, Feb 29, 2016 at 05:22:36PM +0100, Martin Jambor wrote:
> 2016-03-04  Martin Jambor  <mjambor@suse.cz>
> 
> 	* testsuite/lib/libgomp.exp (libgomp_init): Append -Wno-hsa to
> 	ALWAYS_CFLAGS.

Ok.

	Jakub

Re: [hsa testsuite 2/5] Suppress hsa warnings in compiler gomp tests

[Jakub Jelinek]

On Mon, Feb 29, 2016 at 05:02:34PM +0100, Martin Jambor wrote:
> 
> [1] https://gcc.gnu.org/ml/gcc-patches/2016-03/msg00381.html
> 
> 
> 2016-03-04  Martin Jambor  <mjambor@suse.cz>
> 
> 	* c-c++-common/gomp/clauses-1.c: Remove dg-options.
> 	* c-c++-common/gomp/if-1.c: Likewise.
> 	* c-c++-common/gomp/pr61486-2.c: Likewise.
> 	* c-c++-common/gomp/target-teams-1.c: Moved dg-options except -fopenmp
> 	to dg-additional-options.
> 	* g++.dg/gomp/gomp.exp: Pass -Wno-hsa to all tests.
> 	* g++/gomp/target-teams-1.c: Likewise.
> 	* gcc.dg/gomp/gomp.exp: Likewise.
> 	* gcc.dg/gomp/pr68128-2.c: Moved dg-options except -fopenmp to
> 	dg-additional-options.
> 	* gfortran.dg/gomp/gomp.exp: Likewise.
> 	* gfortran.dg/gomp/target1.f90: Remove dg-options.
> 	* gfortran.dg/gomp/target2.f90: Moved dg-options except -fopenmp to
> 	dg-additional-options.
> 	* gfortran.dg/gomp/target3.f90: Remove dg-options.

Ok.

	Jakub

Re: [hsa testsuite 4/5] Adjust libgomp tests that do not work on host fallback

[Jakub Jelinek]

On Mon, Feb 29, 2016 at 05:30:05PM +0100, Martin Jambor wrote:
> Tested both with and without HSA (enabled or present).  OK for trunk?

Ok.  These tests are from the Examples 4.0.x document, so we don't want to
diverge too much on them.

> 2016-02-12  Martin Jambor  <mjambor@suse.cz>
> 
> libgomp/
> 	* testsuite/libgomp.c/examples-4/async_target-2.c: Only run on
> 	non-shared memory accelerators.
> 	* testsuite/libgomp.c/examples-4/device-1.c: Likewise.
> 	* testsuite/libgomp.c/examples-4/target-5.c: Likewise.
> 	* testsuite/libgomp.c/examples-4/target_data-6.c: Likewise.
> 	* testsuite/libgomp.c/examples-4/target_data-7.c: Likewise.
> 	* testsuite/libgomp.fortran/examples-4/async_target-2.f90: Likewise.
> 	* testsuite/libgomp.fortran/examples-4/device-1.f90: Likewise.
> 	* testsuite/libgomp.fortran/examples-4/target-5.f90: Likewise.
> 	* testsuite/libgomp.fortran/examples-4/target_data-6.f90: Likewise.
> 	* testsuite/libgomp.fortran/examples-4/target_data-7.f90: Likewise.

	Jakub

Re: [hsa testsuite 5/5] New directory for HSA-specific C testcases

[Jakub Jelinek]

On Mon, Feb 29, 2016 at 05:31:53PM +0100, Martin Jambor wrote:
> Is the patch OK for trunk?

Ok, thanks.
> 2016-03-03  Martin Jambor  <mjambor@suse.cz>
> 
> 	* testsuite/lib/libgomp.exp
> 	(check_effective_target_hsa_offloading_selected_nocache): New.
> 	(check_effective_target_hsa_offloading_selected): Likewise.
> 	* testsuite/libgomp.hsa.c/c.exp: Likewise.
> 	* testsuite/libgomp.hsa.c/alloca-1.c: Likewise.
> 	* testsuite/libgomp.hsa.c/bitfield-1.c: Likewise.
> 	* testsuite/libgomp.hsa.c/builtins-1.c: Likewise.
> 	* testsuite/libgomp.hsa.c/complex-1.c: Likewise.
> 	* testsuite/libgomp.hsa.c/formal-actual-args-1.c: Likewise.
> 	* testsuite/libgomp.hsa.c/function-call-1.c: Likewise.
> 	* testsuite/libgomp.hsa.c/get-level-1.c: Likewise.
> 	* testsuite/libgomp.hsa.c/gridify-1.c: Likewise.
> 	* testsuite/libgomp.hsa.c/gridify-2.c: Likewise.
> 	* testsuite/libgomp.hsa.c/gridify-3.c: Likewise.
> 	* testsuite/libgomp.hsa.c/gridify-4.c: Likewise.
> 	* testsuite/libgomp.hsa.c/memory-operations-1.c: Likewise.
> 	* testsuite/libgomp.hsa.c/pr69568.c: Likewise.
> 	* testsuite/libgomp.hsa.c/rotate-1.c: Likewise.
> 	* testsuite/libgomp.hsa.c/switch-1.c: Likewise.
> 	* testsuite/libgomp.hsa.c/switch-branch-1.c: Likewise.

	Jakub