This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

[gomp4] various backports from trunk


I've applied this patch which backports my recent trunk changes to
gomp-4_0-branch. Specifically, this patch contains

 * nvptx vector state propagation fix, which includes the updated test
   fix for pr70009

 * combined loop clauses fix

Cesar
2016-03-11  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/c/
	* c-parser.c (c_parser_oacc_loop): Update cclauses and clauses
	when calling c_finish_omp_clauses.

	gcc/
	* config/nvptx/nvptx.c (nvptx_gen_shuffle): Add support for QImode
	and HImode registers.

	gcc/cp/
	* parser.c (cp_parser_oacc_loop): Update cclauses and clauses
	when calling c_finish_omp_clauses.

	gcc/testsuite/
	* c-c++-common/goacc/combined-directives-2.c: New test.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/vprop.c: New test.


diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index bbbe26b..5e5f60d 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -13960,9 +13960,9 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
     {
       clauses = c_oacc_split_loop_clauses (clauses, cclauses);
       if (*cclauses)
-	c_finish_omp_clauses (*cclauses, true, false);
+	*cclauses = c_finish_omp_clauses (*cclauses, true, false);
       if (clauses)
-	c_finish_omp_clauses (clauses, true, false);
+	clauses = c_finish_omp_clauses (clauses, true, false);
     }
 
   tree block = c_begin_compound_stmt (true);
diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 492ebd1..5f10a65 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -1301,6 +1301,20 @@ nvptx_gen_shuffle (rtx dst, rtx src, rtx idx, nvptx_shuffle_kind kind)
 	end_sequence ();
       }
       break;
+    case QImode:
+    case HImode:
+      {
+	rtx tmp = gen_reg_rtx (SImode);
+
+	start_sequence ();
+	emit_insn (gen_rtx_SET (tmp, gen_rtx_fmt_e (ZERO_EXTEND, SImode, src)));
+	emit_insn (nvptx_gen_shuffle (tmp, tmp, idx, kind));
+	emit_insn (gen_rtx_SET (dst, gen_rtx_fmt_e (TRUNCATE, GET_MODE (dst),
+						    tmp)));
+	res = get_insns ();
+	end_sequence ();
+      }
+      break;
       
     default:
       gcc_unreachable ();
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index c90e270..9d70ff7 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -35482,9 +35482,9 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
     {
       clauses = c_oacc_split_loop_clauses (clauses, cclauses);
       if (*cclauses)
-	finish_omp_clauses (*cclauses, true, true);
+	*cclauses = finish_omp_clauses (*cclauses, true, true);
       if (clauses)
-	finish_omp_clauses (clauses, true, true);
+	clauses = finish_omp_clauses (clauses, true, true);
     }
 
   tree block = begin_omp_structured_block ();
diff --git a/gcc/testsuite/c-c++-common/goacc/combined-directives-2.c b/gcc/testsuite/c-c++-common/goacc/combined-directives-2.c
new file mode 100644
index 0000000..c51e2f9
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/combined-directives-2.c
@@ -0,0 +1,14 @@
+/* Ensure that bogus clauses aren't propagated in combined loop
+   constructs.  */
+
+int
+main ()
+{
+  int a, i;
+
+#pragma acc parallel loop vector copy(a[0:100]) reduction(+:a) /* { dg-error "'a' does not have pointer or array type" } */
+  for (i = 0; i < 100; i++)
+    a++;
+
+  return a;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c
new file mode 100644
index 0000000..17b9568
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c
@@ -0,0 +1,34 @@
+#include <assert.h>
+
+#define test(type)				\
+void						\
+test_##type ()					\
+{						\
+  signed type b[100];				\
+  signed type i, j, x = -1, y = -1;		\
+						\
+  _Pragma("acc parallel loop copyout (b)")	\
+  for (j = 0; j > -5; j--)			\
+    {						\
+      type c = x+y;                             \
+      _Pragma("acc loop vector")		\
+      for (i = 0; i < 20; i++)			\
+	b[-j*20 + i] = c;			\
+      b[5-j] = c;                               \
+    }						\
+						\
+  for (i = 0; i < 100; i++)			\
+    assert (b[i] == -2);			\
+}
+
+test(char)
+test(short)
+
+int
+main ()
+{
+  test_char ();
+  test_short ();
+
+  return 0;
+}

Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]