@@ -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
+ }]
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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
new file mode 100644
@@ -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);
+}
new file mode 100644
@@ -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);
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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]);
+ }
+}
new file mode 100644
@@ -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;
+}
+
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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));
+}
new file mode 100644
@@ -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;
+}