* [hsa testsuite 5/5] New directory for HSA-specific C testcases
2016-03-07 17:34 [hsa testsuite 0/5] Re-post of all pending patches adjusting testsuite for HSA Martin Jambor
@ 2016-03-07 17:34 ` Martin Jambor
2016-03-07 18:00 ` Jakub Jelinek
2016-03-07 17:34 ` [hsa testsuite 4/5] Adjust libgomp tests that do not work on host fallback Martin Jambor
` (3 subsequent siblings)
4 siblings, 1 reply; 11+ messages in thread
From: Martin Jambor @ 2016-03-07 17:34 UTC (permalink / raw)
To: GCC Patches; +Cc: Mike Stump
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
^ permalink raw reply [flat|nested] 11+ messages in thread* [hsa testsuite 2/5] Suppress hsa warnings in compiler gomp tests
2016-03-07 17:34 [hsa testsuite 0/5] Re-post of all pending patches adjusting testsuite for HSA Martin Jambor
2016-03-07 17:34 ` [hsa testsuite 5/5] New directory for HSA-specific C testcases Martin Jambor
2016-03-07 17:34 ` [hsa testsuite 4/5] Adjust libgomp tests that do not work on host fallback Martin Jambor
@ 2016-03-07 17:34 ` Martin Jambor
2016-03-07 17:57 ` Jakub Jelinek
2016-03-07 17:34 ` [hsa testsuite 1/5] Gridification tests Martin Jambor
2016-03-07 17:34 ` [hsa testsuite 3/5] Suppress hsa warnings in libgomp tests Martin Jambor
4 siblings, 1 reply; 11+ messages in thread
From: Martin Jambor @ 2016-03-07 17:34 UTC (permalink / raw)
To: GCC Patches; +Cc: Jakub Jelinek
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
^ permalink raw reply [flat|nested] 11+ messages in thread