Merged current trunk to branch.

12 files changed, 189 insertions, 10 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 4c9f9d1..574a9f7 100644
--- a/libgomp/testsuite/libgomp.c/affinity-1.c
+++ b/libgomp/testsuite/libgomp.c/affinity-1.c
@@ -1,5 +1,5 @@
 /* Affinity tests.
-   Copyright (C) 2013-2020 Free Software Foundation, Inc.
+   Copyright (C) 2013-2021 Free Software Foundation, Inc.
 
    GCC is free software; you can redistribute it and/or modify it under
    the terms of the GNU General Public License as published by the Free
@@ -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/sort-1.c b/libgomp/testsuite/libgomp.c/sort-1.c
index bcabea5..eb7f419 100644
--- a/libgomp/testsuite/libgomp.c/sort-1.c
+++ b/libgomp/testsuite/libgomp.c/sort-1.c
@@ -1,5 +1,5 @@
 /* Test and benchmark of a couple of parallel sorting algorithms.
-   Copyright (C) 2008-2020 Free Software Foundation, Inc.
+   Copyright (C) 2008-2021 Free Software Foundation, Inc.
 
    GCC is free software; you can redistribute it and/or modify it under
    the terms of the GNU General Public License as published by the Free
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;
+}