aboutsummaryrefslogtreecommitdiff
path: root/libgomp/testsuite/libgomp.c
diff options
context:
space:
mode:
Diffstat (limited to 'libgomp/testsuite/libgomp.c')
-rw-r--r--libgomp/testsuite/libgomp.c/address-space-1.c28
-rw-r--r--libgomp/testsuite/libgomp.c/affinity-1.c14
-rw-r--r--libgomp/testsuite/libgomp.c/omp-nested-3.c1
-rw-r--r--libgomp/testsuite/libgomp.c/pr46032-2.c1
-rw-r--r--libgomp/testsuite/libgomp.c/pr81778.c48
-rw-r--r--libgomp/testsuite/libgomp.c/pr86416-1.c4
-rw-r--r--libgomp/testsuite/libgomp.c/pr86416-2.c4
-rw-r--r--libgomp/testsuite/libgomp.c/pr99555-1.c21
-rw-r--r--libgomp/testsuite/libgomp.c/target-43.c26
-rw-r--r--libgomp/testsuite/libgomp.c/target-44.c27
-rw-r--r--libgomp/testsuite/libgomp.c/task-reduction-4.c21
11 files changed, 187 insertions, 8 deletions
diff --git a/libgomp/testsuite/libgomp.c/address-space-1.c b/libgomp/testsuite/libgomp.c/address-space-1.c
new file mode 100644
index 0000000..6ad57de
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/address-space-1.c
@@ -0,0 +1,28 @@
+/* Verify OMP instances of variables with address space. */
+
+/* { dg-do run { target i?86-*-* x86_64-*-* } } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* With Intel MIC (emulated) offloading:
+ offload error: process on the device 0 unexpectedly exited with code 0
+ { dg-xfail-run-if TODO { offload_device_intel_mic } } */
+
+#include <assert.h>
+
+int __seg_fs a;
+
+int
+main (void)
+{
+ // a = 123; // SIGSEGV
+ int b;
+#pragma omp target map(alloc: a) map(from: b)
+ {
+ a = 321; // no SIGSEGV (given 'offload_device_nonshared_as')
+ asm volatile ("" : : "g" (&a) : "memory");
+ b = a;
+ }
+ assert (b == 321);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/affinity-1.c b/libgomp/testsuite/libgomp.c/affinity-1.c
index 13a743e..574a9f7 100644
--- a/libgomp/testsuite/libgomp.c/affinity-1.c
+++ b/libgomp/testsuite/libgomp.c/affinity-1.c
@@ -183,20 +183,26 @@ main ()
int test_false = env_proc_bind && strcmp (env_proc_bind, "false") == 0;
int test_true = env_proc_bind && strcmp (env_proc_bind, "true") == 0;
int test_spread_master_close
- = env_proc_bind && strcmp (env_proc_bind, "spread,master,close") == 0;
+ = (env_proc_bind
+ && (strcmp (env_proc_bind, "spread,master,close") == 0
+ || strcmp (env_proc_bind, "spread,primary,close") == 0));
char *env_places = getenv ("OMP_PLACES");
int test_places = 0;
+ if (omp_proc_bind_master != omp_proc_bind_primary)
+ abort ();
+
#ifdef DO_FORK
if (env_places == NULL && contig_cpucount >= 8 && test_false
&& getenv ("GOMP_AFFINITY") == NULL)
{
int i, j, status;
pid_t pid;
- for (j = 0; j < 2; j++)
+ for (j = 0; j < 3; j++)
{
- if (setenv ("OMP_PROC_BIND", j ? "spread,master,close" : "true", 1)
- < 0)
+ if (setenv ("OMP_PROC_BIND",
+ j > 1 ? "spread,primary,close"
+ : (j ? "spread,master,close" : "true"), 1) < 0)
break;
for (i = sizeof (places_array) / sizeof (places_array[0]) - 1;
i; --i)
diff --git a/libgomp/testsuite/libgomp.c/omp-nested-3.c b/libgomp/testsuite/libgomp.c/omp-nested-3.c
index 7790c58..446e6bd 100644
--- a/libgomp/testsuite/libgomp.c/omp-nested-3.c
+++ b/libgomp/testsuite/libgomp.c/omp-nested-3.c
@@ -1,4 +1,5 @@
// { dg-do run { target lto } }
// { dg-additional-options "-fipa-pta -flto -flto-partition=max" }
+// { dg-prune-output "warning: using serial compilation" }
#include "omp-nested-1.c"
diff --git a/libgomp/testsuite/libgomp.c/pr46032-2.c b/libgomp/testsuite/libgomp.c/pr46032-2.c
index 1125f6e..36f3730 100644
--- a/libgomp/testsuite/libgomp.c/pr46032-2.c
+++ b/libgomp/testsuite/libgomp.c/pr46032-2.c
@@ -1,4 +1,5 @@
/* { dg-do run { target lto } } */
/* { dg-options "-O2 -ftree-vectorize -std=c99 -fipa-pta -flto -flto-partition=max" } */
+/* { dg-prune-output "warning: using serial compilation" } */
#include "pr46032.c"
diff --git a/libgomp/testsuite/libgomp.c/pr81778.c b/libgomp/testsuite/libgomp.c/pr81778.c
new file mode 100644
index 0000000..571668e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr81778.c
@@ -0,0 +1,48 @@
+/* Minimized from for-5.c. */
+
+#include <stdio.h>
+#include <stdlib.h>
+
+/* Size of array we want to write. */
+#define N 32
+
+/* Size of extra space before and after. */
+#define CANARY_SIZE (N * 32)
+
+/* Start of array we want to write. */
+#define BASE (CANARY_SIZE)
+
+// Total size to be allocated.
+#define ALLOC_SIZE (CANARY_SIZE + N + CANARY_SIZE)
+
+#pragma omp declare target
+int a[ALLOC_SIZE];
+#pragma omp end declare target
+
+int
+main (void)
+{
+ /* Use variable step in for loop. */
+ int s = 1;
+
+#pragma omp target update to(a)
+
+ /* Write a[BASE] .. a[BASE + N - 1]. */
+#pragma omp target simd
+ for (int i = N - 1; i > -1; i -= s)
+ a[BASE + i] = 1;
+
+#pragma omp target update from(a)
+
+ for (int i = 0; i < ALLOC_SIZE; i++)
+ {
+ int expected = (BASE <= i && i < BASE + N) ? 1 : 0;
+ if (a[i] == expected)
+ continue;
+
+ printf ("Expected %d, got %d at base[%d]\n", expected, a[i], i - BASE);
+ abort ();
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/pr86416-1.c b/libgomp/testsuite/libgomp.c/pr86416-1.c
index ad9370f..6d38692 100644
--- a/libgomp/testsuite/libgomp.c/pr86416-1.c
+++ b/libgomp/testsuite/libgomp.c/pr86416-1.c
@@ -2,8 +2,8 @@
/* { dg-require-effective-target large_long_double } */
/* PR middle-end/86416 */
-/* { dg-error "bit-precision floating-point numbers unsupported .mode '.F'." "" { target offload_device } 0 } */
-/* { dg-excess-errors "Follow-up errors from mkoffload and lto-wrapper" { target offload_device } } */
+/* { dg-error "bit-precision floating-point numbers unsupported .mode '.F'." "" { target { offload_target_nvptx || offload_target_amdgcn } } 0 } */
+/* { dg-excess-errors "Follow-up errors from mkoffload and lto-wrapper" { target { offload_target_nvptx || offload_target_amdgcn } } } */
#include <stdlib.h> /* For abort. */
diff --git a/libgomp/testsuite/libgomp.c/pr86416-2.c b/libgomp/testsuite/libgomp.c/pr86416-2.c
index ec45e40..cffeb3f 100644
--- a/libgomp/testsuite/libgomp.c/pr86416-2.c
+++ b/libgomp/testsuite/libgomp.c/pr86416-2.c
@@ -2,8 +2,8 @@
/* { dg-add-options __float128 } */
/* PR middle-end/86416 */
-/* { dg-error "bit-precision floating-point numbers unsupported .mode '.F'." "" { target offload_device } 0 } */
-/* { dg-excess-errors "Follow-up errors from mkoffload and lto-wrapper" { target offload_device } } */
+/* { dg-error "bit-precision floating-point numbers unsupported .mode '.F'." "" { target { offload_target_nvptx || offload_target_amdgcn } } 0 } */
+/* { dg-excess-errors "Follow-up errors from mkoffload and lto-wrapper" { target { offload_target_nvptx || offload_target_amdgcn } } } */
#include <stdlib.h> /* For abort. */
diff --git a/libgomp/testsuite/libgomp.c/pr99555-1.c b/libgomp/testsuite/libgomp.c/pr99555-1.c
new file mode 100644
index 0000000..bd33b93
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr99555-1.c
@@ -0,0 +1,21 @@
+// PR99555 "[OpenMP/nvptx] Execution-time hang for simple nested OpenMP 'target'/'parallel'/'task' constructs"
+
+// { dg-additional-options "-O0" }
+
+#include <unistd.h> // For 'alarm'.
+
+#include "../libgomp.c-c++-common/on_device_arch.h"
+
+int main (void)
+{
+ if (on_device_arch_nvptx ())
+ alarm (4); /*TODO Until resolved, make sure that we exit quickly, with error status.
+ { dg-xfail-run-if "PR99555" { offload_device_nvptx } } */
+
+#pragma omp target
+#pragma omp parallel // num_threads(1)
+#pragma omp task
+ ;
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-43.c b/libgomp/testsuite/libgomp.c/target-43.c
new file mode 100644
index 0000000..028e912
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-43.c
@@ -0,0 +1,26 @@
+/* { dg-do run } */
+/* { dg-additional-options "-foffload-options=nvptx-none=-latomic" { target { offload_target_nvptx } } } */
+
+#include <stdlib.h>
+
+#define N 32
+#define TYPE char
+
+int
+main (void)
+{
+ TYPE result = 1;
+ TYPE a[N];
+ for (int x = 0; x < N; ++x)
+ a[x] = 1;
+
+#pragma omp target map(tofrom: result) map(to:a)
+#pragma omp for simd reduction(&&:result)
+ for (int x = 0; x < N; ++x)
+ result = result && a[x];
+
+ if (result != 1)
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-44.c b/libgomp/testsuite/libgomp.c/target-44.c
new file mode 100644
index 0000000..a5da81d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-44.c
@@ -0,0 +1,27 @@
+/* { dg-additional-options "-foffload-options=nvptx-none=-latomic" { target { offload_target_nvptx } } } */
+
+#include <stdlib.h>
+
+struct s
+{
+ int i;
+};
+
+#pragma omp declare reduction(+: struct s: omp_out.i += omp_in.i)
+
+int
+main (void)
+{
+ const int N0 = 32768;
+
+ struct s counter_N0 = { 0 };
+#pragma omp target
+#pragma omp for simd reduction(+: counter_N0)
+ for (int i0 = 0 ; i0 < N0 ; i0++ )
+ counter_N0.i += 1;
+
+ if (counter_N0.i != N0)
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/task-reduction-4.c b/libgomp/testsuite/libgomp.c/task-reduction-4.c
new file mode 100644
index 0000000..7ca1d02
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/task-reduction-4.c
@@ -0,0 +1,21 @@
+/* PR middle-end/100471 */
+
+extern void abort (void);
+
+int c;
+
+int
+main ()
+{
+#pragma omp parallel
+#pragma omp single
+ {
+ int r = 0, i;
+ #pragma omp taskloop reduction(+:r)
+ for (i = 0; i < c; i++)
+ r++;
+ if (r != 0)
+ abort ();
+ }
+ return 0;
+}
generated by cgit v1.1 at 2025-09-27 13:59:44 +0000