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

2016-03-07  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.

From-SVN: r234047

libgomp/ChangeLog

 2016-03-07  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.
+
+2016-03-07  Martin Jambor  <mjambor@suse.cz>
+
 	* 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.

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
+    }]
+}

libgomp/testsuite/libgomp.hsa.c/alloca-1.c

+#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;
+}

libgomp/testsuite/libgomp.hsa.c/bitfield-1.c

+#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;
+}

libgomp/testsuite/libgomp.hsa.c/builtins-1.c

+/* { 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;
+}

libgomp/testsuite/libgomp.hsa.c/c.exp

+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

libgomp/testsuite/libgomp.hsa.c/complex-1.c

+#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);
+}

libgomp/testsuite/libgomp.hsa.c/formal-actual-args-1.c

+#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);
+}

libgomp/testsuite/libgomp.hsa.c/function-call-1.c

+#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;
+}

libgomp/testsuite/libgomp.hsa.c/get-level-1.c

+#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;
+}

libgomp/testsuite/libgomp.hsa.c/gridify-1.c

+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;
+}

libgomp/testsuite/libgomp.hsa.c/gridify-2.c

+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;
+}

libgomp/testsuite/libgomp.hsa.c/gridify-3.c

+#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;
+}

libgomp/testsuite/libgomp.hsa.c/gridify-4.c

+#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;
+}

libgomp/testsuite/libgomp.hsa.c/memory-operations-1.c

+#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]);
+    }
+}

libgomp/testsuite/libgomp.hsa.c/pr69568.c

+/* 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;
+}
+

libgomp/testsuite/libgomp.hsa.c/rotate-1.c

+#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;
+}

libgomp/testsuite/libgomp.hsa.c/switch-1.c

+#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));
+}

libgomp/testsuite/libgomp.hsa.c/switch-branch-1.c

+#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;
+}