tree-core.h (enum omp_clause_code): Add OMP_CLAUSE__SCANTEMP_ clause.

	* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE__SCANTEMP_
	clause.
	* tree.h (OMP_CLAUSE_DECL): Use OMP_CLAUSE__SCANTEMP_ instead of
	OMP_CLAUSE__CONDTEMP_ as range's upper bound.
	(OMP_CLAUSE__SCANTEMP__ALLOC, OMP_CLAUSE__SCANTEMP__CONTROL): Define.
	* tree.c (omp_clause_num_ops, omp_clause_code_name): Add
	OMP_CLAUSE__SCANTEMP_ entry.
	(walk_tree_1): Handle OMP_CLAUSE__SCANTEMP_.
	* tree-pretty-print.c (dump_omp_clause): Likewise.
	* tree-nested.c (convert_nonlocal_omp_clauses,
	convert_local_omp_clauses): Likewise.
	* omp-general.h (struct omp_for_data): Add have_scantemp and
	have_nonctrl_scantemp members.
	* omp-general.c (omp_extract_for_data): Initialize them.
	* omp-low.c (struct omp_context): Add scan_exclusive member.
	(scan_omp_1_stmt): Don't unnecessarily mask gimple_omp_for_kind
	result again with GF_OMP_FOR_KIND_MASK.  Initialize also
	ctx->scan_exclusive.
	(lower_rec_simd_input_clauses): Use ctx->scan_exclusive instead
	of !ctx->scan_inclusive.
	(lower_rec_input_clauses): Simplify gimplification of dtors using
	gimplify_and_add.  For non-is_simd test OMP_CLAUSE_REDUCTION_INSCAN
	rather than rvarp.  Handle OMP_CLAUSE_REDUCTION_INSCAN in worksharing
	loops.  Don't add barrier for reduction_omp_orig_ref if
	ctx->scan_??xclusive.
	(lower_reduction_clauses): Don't do anything for ctx->scan_??xclusive.
	(lower_omp_scan): Use ctx->scan_exclusive instead
	of !ctx->scan_inclusive.  Handle worksharing loops with inscan
	reductions.  Use new_vard != new_var instead of repeated
	omp_is_reference calls.
	(omp_find_scan, lower_omp_for_scan): New functions.
	(lower_omp_for): Call lower_omp_for_scan for worksharing loops with
	inscan reductions.
	* omp-expand.c (expand_omp_scantemp_alloc): New function.
	(expand_omp_for_static_nochunk): Handle fd->have_nonctrl_scantemp
	and fd->have_scantemp.

	* c-c++-common/gomp/scan-3.c (f1): Don't expect a sorry message.
	* c-c++-common/gomp/scan-5.c (foo): Likewise.

	* testsuite/libgomp.c++/scan-1.C: New test.
	* testsuite/libgomp.c++/scan-2.C: New test.
	* testsuite/libgomp.c++/scan-3.C: New test.
	* testsuite/libgomp.c++/scan-4.C: New test.
	* testsuite/libgomp.c++/scan-5.C: New test.
	* testsuite/libgomp.c++/scan-6.C: New test.
	* testsuite/libgomp.c++/scan-7.C: New test.
	* testsuite/libgomp.c++/scan-8.C: New test.
	* testsuite/libgomp.c/scan-1.c: New test.
	* testsuite/libgomp.c/scan-2.c: New test.
	* testsuite/libgomp.c/scan-3.c: New test.
	* testsuite/libgomp.c/scan-4.c: New test.
	* testsuite/libgomp.c/scan-5.c: New test.
	* testsuite/libgomp.c/scan-6.c: New test.
	* testsuite/libgomp.c/scan-7.c: New test.
	* testsuite/libgomp.c/scan-8.c: New test.

From-SVN: r272958

8 files changed, 996 insertions, 0 deletions

diff --git a/libgomp/testsuite/libgomp.c/scan-1.c b/libgomp/testsuite/libgomp.c/scan-1.c
new file mode 100644
index 0000000..dd34f33
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/scan-1.c
@@ -0,0 +1,115 @@
+/* { dg-require-effective-target size32plus } */
+
+extern void abort (void);
+int r, a[1024], b[1024];
+
+__attribute__((noipa)) void
+foo (int *a, int *b)
+{
+  #pragma omp for reduction (inscan, +:r)
+  for (int i = 0; i < 1024; i++)
+    {
+      r += a[i];
+      #pragma omp scan inclusive(r)
+      b[i] = r;
+    }
+}
+
+__attribute__((noipa)) int
+bar (void)
+{
+  int s = 0;
+  #pragma omp parallel
+  #pragma omp for reduction (inscan, +:s)
+  for (int i = 0; i < 1024; i++)
+    {
+      s += 2 * a[i];
+      #pragma omp scan inclusive(s)
+      b[i] = s;
+    }
+  return s;
+}
+
+__attribute__((noipa)) void
+baz (int *a, int *b)
+{
+  #pragma omp parallel for reduction (inscan, +:r)
+  for (int i = 0; i < 1024; i++)
+    {
+      r += a[i];
+      #pragma omp scan inclusive(r)
+      b[i] = r;
+    }
+}
+
+__attribute__((noipa)) int
+qux (void)
+{
+  int s = 0;
+  #pragma omp parallel for reduction (inscan, +:s)
+  for (int i = 0; i < 1024; i++)
+    {
+      s += 2 * a[i];
+      #pragma omp scan inclusive(s)
+      b[i] = s;
+    }
+  return s;
+}
+
+int
+main ()
+{
+  int s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      a[i] = i;
+      b[i] = -1;
+      asm ("" : "+g" (i));
+    }
+  #pragma omp parallel
+  foo (a, b);
+  if (r != 1024 * 1023 / 2)
+    abort ();
+  for (int i = 0; i < 1024; ++i)
+    {
+      s += i;
+      if (b[i] != s)
+	abort ();
+      else
+	b[i] = 25;
+    }
+  if (bar () != 1024 * 1023)
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      s += 2 * i;
+      if (b[i] != s)
+	abort ();
+      else
+	b[i] = -1;
+    }
+  r = 0;
+  baz (a, b);
+  if (r != 1024 * 1023 / 2)
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      s += i;
+      if (b[i] != s)
+	abort ();
+      else
+	b[i] = -25;
+    }
+  if (qux () != 1024 * 1023)
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      s += 2 * i;
+      if (b[i] != s)
+	abort ();
+    }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/scan-2.c b/libgomp/testsuite/libgomp.c/scan-2.c
new file mode 100644
index 0000000..476bbed
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/scan-2.c
@@ -0,0 +1,117 @@
+/* { dg-require-effective-target size32plus } */
+
+extern void abort (void);
+int r, a[1024], b[1024];
+
+#pragma omp declare reduction (foo: int: omp_out += omp_in) initializer (omp_priv = 0)
+
+__attribute__((noipa)) void
+foo (int *a, int *b)
+{
+  #pragma omp for reduction (inscan, foo:r)
+  for (int i = 0; i < 1024; i++)
+    {
+      r += a[i];
+      #pragma omp scan inclusive(r)
+      b[i] = r;
+    }
+}
+
+__attribute__((noipa)) int
+bar (void)
+{
+  int s = 0;
+  #pragma omp parallel
+  #pragma omp for reduction (inscan, foo:s)
+  for (int i = 0; i < 1024; i++)
+    {
+      s += 2 * a[i];
+      #pragma omp scan inclusive(s)
+      b[i] = s;
+    }
+  return s;
+}
+
+__attribute__((noipa)) void
+baz (int *a, int *b)
+{
+  #pragma omp parallel for reduction (inscan, foo:r)
+  for (int i = 0; i < 1024; i++)
+    {
+      r += a[i];
+      #pragma omp scan inclusive(r)
+      b[i] = r;
+    }
+}
+
+__attribute__((noipa)) int
+qux (void)
+{
+  int s = 0;
+  #pragma omp parallel for reduction (inscan, foo:s)
+  for (int i = 0; i < 1024; i++)
+    {
+      s += 2 * a[i];
+      #pragma omp scan inclusive(s)
+      b[i] = s;
+    }
+  return s;
+}
+
+int
+main ()
+{
+  int s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      a[i] = i;
+      b[i] = -1;
+      asm ("" : "+g" (i));
+    }
+  #pragma omp parallel
+  foo (a, b);
+  if (r != 1024 * 1023 / 2)
+    abort ();
+  for (int i = 0; i < 1024; ++i)
+    {
+      s += i;
+      if (b[i] != s)
+	abort ();
+      else
+	b[i] = 25;
+    }
+  if (bar () != 1024 * 1023)
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      s += 2 * i;
+      if (b[i] != s)
+	abort ();
+      else
+	b[i] = -1;
+    }
+  r = 0;
+  baz (a, b);
+  if (r != 1024 * 1023 / 2)
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      s += i;
+      if (b[i] != s)
+	abort ();
+      else
+	b[i] = -25;
+    }
+  if (qux () != 1024 * 1023)
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      s += 2 * i;
+      if (b[i] != s)
+	abort ();
+    }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/scan-3.c b/libgomp/testsuite/libgomp.c/scan-3.c
new file mode 100644
index 0000000..2e1fcad
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/scan-3.c
@@ -0,0 +1,88 @@
+/* { dg-require-effective-target size32plus } */
+
+extern void abort (void);
+float r = 1.0f, a[1024], b[1024];
+
+__attribute__((noipa)) void
+foo (float *a, float *b)
+{
+  #pragma omp for reduction (inscan, *:r)
+  for (int i = 0; i < 1024; i++)
+    {
+      r *= a[i];
+      #pragma omp scan inclusive(r)
+      b[i] = r;
+    }
+}
+
+__attribute__((noipa)) float
+bar (void)
+{
+  float s = -__builtin_inff ();
+  #pragma omp parallel for reduction (inscan, max:s)
+  for (int i = 0; i < 1024; i++)
+    {
+      s = s > a[i] ? s : a[i];
+      #pragma omp scan inclusive(s)
+      b[i] = s;
+    }
+  return s;
+}
+
+int
+main ()
+{
+  float s = 1.0f;
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (i < 80)
+	a[i] = (i & 1) ? 0.25f : 0.5f;
+      else if (i < 200)
+	a[i] = (i % 3) == 0 ? 2.0f : (i % 3) == 1 ? 4.0f : 1.0f;
+      else if (i < 280)
+	a[i] = (i & 1) ? 0.25f : 0.5f;
+      else if (i < 380)
+	a[i] = (i % 3) == 0 ? 2.0f : (i % 3) == 1 ? 4.0f : 1.0f;
+      else
+	switch (i % 6)
+	  {
+	  case 0: a[i] = 0.25f; break;
+	  case 1: a[i] = 2.0f; break;
+	  case 2: a[i] = -1.0f; break;
+	  case 3: a[i] = -4.0f; break;
+	  case 4: a[i] = 0.5f; break;
+	  case 5: a[i] = 1.0f; break;
+	  default: a[i] = 0.0f; break;
+	  }
+      b[i] = -19.0f;
+      asm ("" : "+g" (i));
+    }
+  #pragma omp parallel
+  foo (a, b);
+  if (r * 16384.0f != 0.125f)
+    abort ();
+  float m = -175.25f;
+  for (int i = 0; i < 1024; ++i)
+    {
+      s *= a[i];
+      if (b[i] != s)
+	abort ();
+      else
+	{
+	  a[i] = m - ((i % 3) == 1 ? 2.0f : (i % 3) == 2 ? 4.0f : 0.0f);
+	  b[i] = -231.75f;
+	  m += 0.75f;
+	}
+    }
+  if (bar () != 592.0f)
+    abort ();
+  s = -__builtin_inff ();
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (s < a[i])
+	s = a[i];
+      if (b[i] != s)
+	abort ();
+    }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/scan-4.c b/libgomp/testsuite/libgomp.c/scan-4.c
new file mode 100644
index 0000000..b1e9993
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/scan-4.c
@@ -0,0 +1,179 @@
+/* { dg-require-effective-target size32plus } */
+
+extern void abort (void);
+int r, a[1024], b[1024];
+unsigned short r2, b2[1024];
+unsigned char r3, b3[1024];
+
+__attribute__((noipa)) void
+foo (int *a, int *b, unsigned short *b2, unsigned char *b3)
+{
+  #pragma omp for reduction (inscan, +:r, r2, r3)
+  for (int i = 0; i < 1024; i++)
+    {
+      { r += a[i]; r2 += a[i]; r3 += a[i]; }
+      #pragma omp scan inclusive(r, r2, r3)
+      {
+	b[i] = r;
+	b2[i] = r2;
+	b3[i] = r3;
+      }
+    }
+}
+
+__attribute__((noipa)) int
+bar (unsigned short *s2p, unsigned char *s3p)
+{
+  int s = 0;
+  unsigned short s2 = 0;
+  unsigned char s3 = 0;
+  #pragma omp parallel
+  #pragma omp for reduction (inscan, +:s, s2, s3)
+  for (int i = 0; i < 1024; i++)
+    {
+      {
+	s += 2 * a[i];
+	s2 += 2 * a[i];
+	s3 += 2 * a[i];
+      }
+      #pragma omp scan inclusive(s, s2, s3)
+      { b[i] = s; b2[i] = s2; b3[i] = s3; }
+    }
+  *s2p = s2;
+  *s3p = s3;
+  return s;
+}
+
+__attribute__((noipa)) void
+baz (int *a, int *b, unsigned short *b2, unsigned char *b3)
+{
+  #pragma omp parallel for reduction (inscan, +:r, r2, r3)
+  for (int i = 0; i < 1024; i++)
+    {
+      {
+	r += a[i];
+	r2 += a[i];
+	r3 += a[i];
+      }
+      #pragma omp scan inclusive(r, r2, r3)
+      {
+	b[i] = r;
+	b2[i] = r2;
+	b3[i] = r3;
+      }
+    }
+}
+
+__attribute__((noipa)) int
+qux (unsigned short *s2p, unsigned char *s3p)
+{
+  int s = 0;
+  unsigned short s2 = 0;
+  unsigned char s3 = 0;
+  #pragma omp parallel for reduction (inscan, +:s, s2, s3)
+  for (int i = 0; i < 1024; i++)
+    {
+      { s += 2 * a[i]; s2 += 2 * a[i]; s3 += 2 * a[i]; }
+      #pragma omp scan inclusive(s, s2, s3)
+      { b[i] = s; b2[i] = s2; b3[i] = s3; }
+    }
+  *s2p = s2;
+  *s3p = s3;
+  return s;
+}
+
+int
+main ()
+{
+  int s = 0;
+  unsigned short s2;
+  unsigned char s3;
+  for (int i = 0; i < 1024; ++i)
+    {
+      a[i] = i;
+      b[i] = -1;
+      b2[i] = -1;
+      b3[i] = -1;
+      asm ("" : "+g" (i));
+    }
+  #pragma omp parallel
+  foo (a, b, b2, b3);
+  if (r != 1024 * 1023 / 2
+      || r2 != (unsigned short) r
+      || r3 != (unsigned char) r)
+    abort ();
+  for (int i = 0; i < 1024; ++i)
+    {
+      s += i;
+      if (b[i] != s
+	  || b2[i] != (unsigned short) s
+	  || b3[i] != (unsigned char) s)
+	abort ();
+      else
+	{
+	  b[i] = 25;
+	  b2[i] = 24;
+	  b3[i] = 26;
+	}
+    }
+  if (bar (&s2, &s3) != 1024 * 1023)
+    abort ();
+  if (s2 != (unsigned short) (1024 * 1023)
+      || s3 != (unsigned char) (1024 * 1023))
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      s += 2 * i;
+      if (b[i] != s
+	  || b2[i] != (unsigned short) s
+	  || b3[i] != (unsigned char) s)
+	abort ();
+      else
+	{
+	  b[i] = -1;
+	  b2[i] = -1;
+	  b3[i] = -1;
+	}
+    }
+  r = 0;
+  r2 = 0;
+  r3 = 0;
+  baz (a, b, b2, b3);
+  if (r != 1024 * 1023 / 2
+      || r2 != (unsigned short) r
+      || r3 != (unsigned char) r)
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      s += i;
+      if (b[i] != s
+	  || b2[i] != (unsigned short) s
+	  || b3[i] != (unsigned char) s)
+	abort ();
+      else
+	{
+	  b[i] = 25;
+	  b2[i] = 24;
+	  b3[i] = 26;
+	}
+    }
+  s2 = 0;
+  s3 = 0;
+  if (qux (&s2, &s3) != 1024 * 1023)
+    abort ();
+  if (s2 != (unsigned short) (1024 * 1023)
+      || s3 != (unsigned char) (1024 * 1023))
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      s += 2 * i;
+      if (b[i] != s
+	  || b2[i] != (unsigned short) s
+	  || b3[i] != (unsigned char) s)
+	abort ();
+    }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/scan-5.c b/libgomp/testsuite/libgomp.c/scan-5.c
new file mode 100644
index 0000000..c504182
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/scan-5.c
@@ -0,0 +1,115 @@
+/* { dg-require-effective-target size32plus } */
+
+extern void abort (void);
+int r, a[1024], b[1024];
+
+__attribute__((noipa)) void
+foo (int *a, int *b)
+{
+  #pragma omp for reduction (inscan, +:r)
+  for (int i = 0; i < 1024; i++)
+    {
+      b[i] = r;
+      #pragma omp scan exclusive(r)
+      r += a[i];
+    }
+}
+
+__attribute__((noipa)) int
+bar (void)
+{
+  int s = 0;
+  #pragma omp parallel
+  #pragma omp for reduction (inscan, +:s)
+  for (int i = 0; i < 1024; i++)
+    {
+      b[i] = s;
+      #pragma omp scan exclusive(s)
+      s += 2 * a[i];
+    }
+  return s;
+}
+
+__attribute__((noipa)) void
+baz (int *a, int *b)
+{
+  #pragma omp parallel for reduction (inscan, +:r)
+  for (int i = 0; i < 1024; i++)
+    {
+      b[i] = r;
+      #pragma omp scan exclusive(r)
+      r += a[i];
+    }
+}
+
+__attribute__((noipa)) int
+qux (void)
+{
+  int s = 0;
+  #pragma omp parallel for reduction (inscan, +:s)
+  for (int i = 0; i < 1024; i++)
+    {
+      b[i] = s;
+      #pragma omp scan exclusive(s)
+      s += 2 * a[i];
+    }
+  return s;
+}
+
+int
+main ()
+{
+  int s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      a[i] = i;
+      b[i] = -1;
+      asm ("" : "+g" (i));
+    }
+  #pragma omp parallel
+  foo (a, b);
+  if (r != 1024 * 1023 / 2)
+    abort ();
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (b[i] != s)
+	abort ();
+      else
+	b[i] = 25;
+      s += i;
+    }
+  if (bar () != 1024 * 1023)
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (b[i] != s)
+	abort ();
+      else
+	b[i] = -1;
+      s += 2 * i;
+    }
+  r = 0;
+  baz (a, b);
+  if (r != 1024 * 1023 / 2)
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (b[i] != s)
+	abort ();
+      else
+	b[i] = -25;
+      s += i;
+    }
+  if (qux () != 1024 * 1023)
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (b[i] != s)
+	abort ();
+      s += 2 * i;
+    }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/scan-6.c b/libgomp/testsuite/libgomp.c/scan-6.c
new file mode 100644
index 0000000..9434ebb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/scan-6.c
@@ -0,0 +1,117 @@
+/* { dg-require-effective-target size32plus } */
+
+extern void abort (void);
+int r, a[1024], b[1024];
+
+#pragma omp declare reduction (foo: int: omp_out += omp_in) initializer (omp_priv = 0)
+
+__attribute__((noipa)) void
+foo (int *a, int *b)
+{
+  #pragma omp for reduction (inscan, foo:r)
+  for (int i = 0; i < 1024; i++)
+    {
+      b[i] = r;
+      #pragma omp scan exclusive(r)
+      r += a[i];
+    }
+}
+
+__attribute__((noipa)) int
+bar (void)
+{
+  int s = 0;
+  #pragma omp parallel
+  #pragma omp for reduction (inscan, foo:s)
+  for (int i = 0; i < 1024; i++)
+    {
+      b[i] = s;
+      #pragma omp scan exclusive(s)
+      s += 2 * a[i];
+    }
+  return s;
+}
+
+__attribute__((noipa)) void
+baz (int *a, int *b)
+{
+  #pragma omp parallel for reduction (inscan, foo:r)
+  for (int i = 0; i < 1024; i++)
+    {
+      b[i] = r;
+      #pragma omp scan exclusive(r)
+      r += a[i];
+    }
+}
+
+__attribute__((noipa)) int
+qux (void)
+{
+  int s = 0;
+  #pragma omp parallel for reduction (inscan, foo:s)
+  for (int i = 0; i < 1024; i++)
+    {
+      b[i] = s;
+      #pragma omp scan exclusive(s)
+      s += 2 * a[i];
+    }
+  return s;
+}
+
+int
+main ()
+{
+  int s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      a[i] = i;
+      b[i] = -1;
+      asm ("" : "+g" (i));
+    }
+  #pragma omp parallel
+  foo (a, b);
+  if (r != 1024 * 1023 / 2)
+    abort ();
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (b[i] != s)
+	abort ();
+      else
+	b[i] = 25;
+      s += i;
+    }
+  if (bar () != 1024 * 1023)
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (b[i] != s)
+	abort ();
+      else
+	b[i] = -1;
+      s += 2 * i;
+    }
+  r = 0;
+  baz (a, b);
+  if (r != 1024 * 1023 / 2)
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (b[i] != s)
+	abort ();
+      else
+	b[i] = -25;
+      s += i;
+    }
+  if (qux () != 1024 * 1023)
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (b[i] != s)
+	abort ();
+      s += 2 * i;
+    }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/scan-7.c b/libgomp/testsuite/libgomp.c/scan-7.c
new file mode 100644
index 0000000..20b7071
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/scan-7.c
@@ -0,0 +1,86 @@
+/* { dg-require-effective-target size32plus } */
+
+extern void abort (void);
+float r = 1.0f, a[1024], b[1024];
+
+__attribute__((noipa)) void
+foo (float *a, float *b)
+{
+  #pragma omp for reduction (inscan, *:r)
+  for (int i = 0; i < 1024; i++)
+    {
+      b[i] = r;
+      #pragma omp scan exclusive(r)
+      r *= a[i];
+    }
+}
+
+__attribute__((noipa)) float
+bar (void)
+{
+  float s = -__builtin_inff ();
+  #pragma omp parallel for reduction (inscan, max:s)
+  for (int i = 0; i < 1024; i++)
+    {
+      b[i] = s;
+      #pragma omp scan exclusive(s)
+      s = s > a[i] ? s : a[i];
+    }
+  return s;
+}
+
+int
+main ()
+{
+  float s = 1.0f;
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (i < 80)
+	a[i] = (i & 1) ? 0.25f : 0.5f;
+      else if (i < 200)
+	a[i] = (i % 3) == 0 ? 2.0f : (i % 3) == 1 ? 4.0f : 1.0f;
+      else if (i < 280)
+	a[i] = (i & 1) ? 0.25f : 0.5f;
+      else if (i < 380)
+	a[i] = (i % 3) == 0 ? 2.0f : (i % 3) == 1 ? 4.0f : 1.0f;
+      else
+	switch (i % 6)
+	  {
+	  case 0: a[i] = 0.25f; break;
+	  case 1: a[i] = 2.0f; break;
+	  case 2: a[i] = -1.0f; break;
+	  case 3: a[i] = -4.0f; break;
+	  case 4: a[i] = 0.5f; break;
+	  case 5: a[i] = 1.0f; break;
+	  default: a[i] = 0.0f; break;
+	  }
+      b[i] = -19.0f;
+      asm ("" : "+g" (i));
+    }
+  #pragma omp parallel
+  foo (a, b);
+  if (r * 16384.0f != 0.125f)
+    abort ();
+  float m = -175.25f;
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (b[i] != s)
+	abort ();
+      else
+	b[i] = -231.75f;
+      s *= a[i];
+      a[i] = m - ((i % 3) == 1 ? 2.0f : (i % 3) == 2 ? 4.0f : 0.0f);
+      m += 0.75f;
+    }
+  if (bar () != 592.0f)
+    abort ();
+  s = -__builtin_inff ();
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (b[i] != s)
+	abort ();
+      if (s < a[i])
+	s = a[i];
+    }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/scan-8.c b/libgomp/testsuite/libgomp.c/scan-8.c
new file mode 100644
index 0000000..f09c85c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/scan-8.c
@@ -0,0 +1,179 @@
+/* { dg-require-effective-target size32plus } */
+
+extern void abort (void);
+int r, a[1024], b[1024];
+unsigned short r2, b2[1024];
+unsigned char r3, b3[1024];
+
+__attribute__((noipa)) void
+foo (int *a, int *b, unsigned short *b2, unsigned char *b3)
+{
+  #pragma omp for reduction (inscan, +:r, r2, r3)
+  for (int i = 0; i < 1024; i++)
+    {
+      {
+	b[i] = r;
+	b2[i] = r2;
+	b3[i] = r3;
+      }
+      #pragma omp scan exclusive(r, r2, r3)
+      { r += a[i]; r2 += a[i]; r3 += a[i]; }
+    }
+}
+
+__attribute__((noipa)) int
+bar (unsigned short *s2p, unsigned char *s3p)
+{
+  int s = 0;
+  unsigned short s2 = 0;
+  unsigned char s3 = 0;
+  #pragma omp parallel
+  #pragma omp for reduction (inscan, +:s, s2, s3)
+  for (int i = 0; i < 1024; i++)
+    {
+      { b[i] = s; b2[i] = s2; b3[i] = s3; }
+      #pragma omp scan exclusive(s, s2, s3)
+      {
+	s += 2 * a[i];
+	s2 += 2 * a[i];
+	s3 += 2 * a[i];
+      }
+    }
+  *s2p = s2;
+  *s3p = s3;
+  return s;
+}
+
+__attribute__((noipa)) void
+baz (int *a, int *b, unsigned short *b2, unsigned char *b3)
+{
+  #pragma omp parallel for reduction (inscan, +:r, r2, r3)
+  for (int i = 0; i < 1024; i++)
+    {
+      {
+	b[i] = r;
+	b2[i] = r2;
+	b3[i] = r3;
+      }
+      #pragma omp scan exclusive(r, r2, r3)
+      {
+	r += a[i];
+	r2 += a[i];
+	r3 += a[i];
+      }
+    }
+}
+
+__attribute__((noipa)) int
+qux (unsigned short *s2p, unsigned char *s3p)
+{
+  int s = 0;
+  unsigned short s2 = 0;
+  unsigned char s3 = 0;
+  #pragma omp parallel for reduction (inscan, +:s, s2, s3)
+  for (int i = 0; i < 1024; i++)
+    {
+      { b[i] = s; b2[i] = s2; b3[i] = s3; }
+      #pragma omp scan exclusive(s, s2, s3)
+      { s += 2 * a[i]; s2 += 2 * a[i]; s3 += 2 * a[i]; }
+    }
+  *s2p = s2;
+  *s3p = s3;
+  return s;
+}
+
+int
+main ()
+{
+  int s = 0;
+  unsigned short s2;
+  unsigned char s3;
+  for (int i = 0; i < 1024; ++i)
+    {
+      a[i] = i;
+      b[i] = -1;
+      b2[i] = -1;
+      b3[i] = -1;
+      asm ("" : "+g" (i));
+    }
+  #pragma omp parallel
+  foo (a, b, b2, b3);
+  if (r != 1024 * 1023 / 2
+      || r2 != (unsigned short) r
+      || r3 != (unsigned char) r)
+    abort ();
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (b[i] != s
+	  || b2[i] != (unsigned short) s
+	  || b3[i] != (unsigned char) s)
+	abort ();
+      else
+	{
+	  b[i] = 25;
+	  b2[i] = 24;
+	  b3[i] = 26;
+	}
+      s += i;
+    }
+  if (bar (&s2, &s3) != 1024 * 1023)
+    abort ();
+  if (s2 != (unsigned short) (1024 * 1023)
+      || s3 != (unsigned char) (1024 * 1023))
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (b[i] != s
+	  || b2[i] != (unsigned short) s
+	  || b3[i] != (unsigned char) s)
+	abort ();
+      else
+	{
+	  b[i] = -1;
+	  b2[i] = -1;
+	  b3[i] = -1;
+	}
+      s += 2 * i;
+    }
+  r = 0;
+  r2 = 0;
+  r3 = 0;
+  baz (a, b, b2, b3);
+  if (r != 1024 * 1023 / 2
+      || r2 != (unsigned short) r
+      || r3 != (unsigned char) r)
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (b[i] != s
+	  || b2[i] != (unsigned short) s
+	  || b3[i] != (unsigned char) s)
+	abort ();
+      else
+	{
+	  b[i] = 25;
+	  b2[i] = 24;
+	  b3[i] = 26;
+	}
+      s += i;
+    }
+  s2 = 0;
+  s3 = 0;
+  if (qux (&s2, &s3) != 1024 * 1023)
+    abort ();
+  if (s2 != (unsigned short) (1024 * 1023)
+      || s3 != (unsigned char) (1024 * 1023))
+    abort ();
+  s = 0;
+  for (int i = 0; i < 1024; ++i)
+    {
+      if (b[i] != s
+	  || b2[i] != (unsigned short) s
+	  || b3[i] != (unsigned char) s)
+	abort ();
+      s += 2 * i;
+    }
+  return 0;
+}