This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[hsa,testsuite] New directory for HSA-specific C testcases
- From: Martin Jambor <mjambor at suse dot cz>
- To: GCC Patches <gcc-patches at gcc dot gnu dot org>
- Cc: ro at CeBiTec dot Uni-Bielefeld dot DE, mikestump at comcast dot net
- Date: Fri, 26 Feb 2016 17:00:55 +0100
- Subject: [hsa,testsuite] New directory for HSA-specific C testcases
- Authentication-results: sourceware.org; auth=none
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.
The only way to detect that situation I could think of is to run a
simple kernel with environment variable HSA_DEBUG set and look HSA
libgomp plugin debug output in the stderr stream. So I started with a
copy of proc check_runtime_nocache and repurposed it for this task and
then guarded any execution of tests by that predicate in the new c.exp
file. The new directory is in the libgomp testsuite because (at least
now) there is no other way of running hsa stuff but through libgomp.
All tests there are run twice, at -O0 and -O2, because we have found
that both levels are very useful for detecting errors in the lowering
to HSA.
I have very little experience with tcl, expect or DejaGNU and would
appreciate very much any feedback or guidance of anyone more
experience in these areas. In particular, it would probably be better
to structure the new predicate as an effective-target one but I am not
sure what amount of caching would I have to do manually for such a new
runtime test.
After I incorporate all feedback, I would of course like to commit
this to trunk. Needless to say, the patch has been tested and works.
Thanks,
Martin
2016-02-22 Martin Jambor <mjambor@suse.cz>
Martin Liska <mliska@suse.cz>
* testsuite/lib/libgomp.exp (check_hsa_offloading_available): New.
* 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 | 44 ++++++
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 | 41 ++++++
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, 1160 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 154a447..1c917c8 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -392,3 +392,47 @@ proc check_effective_target_openacc_host_selected { } {
}
return 0;
}
+
+# Return 1 if the selected OMP device is actually a HSA device
+
+proc check_hsa_offloading_available {} {
+ 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
+}
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..c746cfb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/c.exp
@@ -0,0 +1,41 @@
+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
+
+if [info exists DEFAULT_CFLAGS] then {
+ set CFLAGS_list [list $DEFAULT_CFLAGS]
+} else {
+ set CFLAGS_list [list "-O0" "-O2"]
+}
+
+if [check_hsa_offloading_available] {
+ foreach USE_CFLAGS $CFLAGS_list {
+ # Gather a list of all tests.
+ set tests [lsort [find $srcdir/$subdir *.c]]
+ # Main loop.
+ dg-runtest $tests "" $USE_CFLAGS
+ }
+}
+
+# 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