From 96a71bd53c4da5f3cb3c72c6f5c7f094ae92caed Mon Sep 17 00:00:00 2001 From: Martin Jambor Date: Mon, 7 Mar 2016 19:45:17 +0100 Subject: [PATCH] [hsa testsuite] New directory for HSA-specific C testcases 2016-03-07 Martin Jambor * 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 | 23 +++ 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 +++++++ .../libgomp.hsa.c/formal-actual-args-1.c | 83 +++++++++ .../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 +++++ .../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 ++++++++++++++++ .../testsuite/libgomp.hsa.c/switch-branch-1.c | 116 +++++++++++++ 19 files changed, 1193 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/ChangeLog b/libgomp/ChangeLog index cfdee56ebeb..ab8cfd47fe3 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,26 @@ +2016-03-07 Martin Jambor + + * 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 * testsuite/libgomp.c/examples-4/async_target-2.c: Only run on diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index bbc2c260831..0d5b6d42d7e 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 00000000000..48dca94a47f --- /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 00000000000..4dbf3481733 --- /dev/null +++ b/libgomp/testsuite/libgomp.hsa.c/bitfield-1.c @@ -0,0 +1,160 @@ +#include + +#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 00000000000..e603c21afcd --- /dev/null +++ b/libgomp/testsuite/libgomp.hsa.c/builtins-1.c @@ -0,0 +1,97 @@ +/* { dg-additional-options "-ffast-math" } */ + +#include +#include + +#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 00000000000..4614192320a --- /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 00000000000..438c64a1593 --- /dev/null +++ b/libgomp/testsuite/libgomp.hsa.c/complex-1.c @@ -0,0 +1,65 @@ +#include +#include +#include + +#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 00000000000..058a036d371 --- /dev/null +++ b/libgomp/testsuite/libgomp.hsa.c/formal-actual-args-1.c @@ -0,0 +1,83 @@ +#include + +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 00000000000..7f15dff96b9 --- /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 00000000000..81c9df00276 --- /dev/null +++ b/libgomp/testsuite/libgomp.hsa.c/get-level-1.c @@ -0,0 +1,26 @@ +#include + +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 00000000000..b670b9b654c --- /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 00000000000..3692eb0d11c --- /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 00000000000..f881d81e18e --- /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 00000000000..c3fbdbf55d4 --- /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 00000000000..a17be932111 --- /dev/null +++ b/libgomp/testsuite/libgomp.hsa.c/memory-operations-1.c @@ -0,0 +1,92 @@ +#include + +#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 00000000000..6262eee3e71 --- /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 00000000000..494388bd10c --- /dev/null +++ b/libgomp/testsuite/libgomp.hsa.c/rotate-1.c @@ -0,0 +1,39 @@ +#include +#include + +#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 00000000000..a180cf6cb7b --- /dev/null +++ b/libgomp/testsuite/libgomp.hsa.c/switch-1.c @@ -0,0 +1,145 @@ +#include + +#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 00000000000..9af1d6d0762 --- /dev/null +++ b/libgomp/testsuite/libgomp.hsa.c/switch-branch-1.c @@ -0,0 +1,116 @@ +#include + +#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; +}