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 <ce...@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; +}