diff options
| author | Tom de Vries <tom@codesourcery.com> | 2018-03-20 10:31:23 +0000 |
|---|---|---|
| committer | Tom de Vries <vries@gcc.gnu.org> | 2018-03-20 10:31:23 +0000 |
| commit | 038012e2d9211e56a2d6710a85314373af9d4d44 (patch) | |
| tree | 47c77648024f770d83b12543cd295c655ee1e2ac /gcc | |
| parent | 452154b9e6d3aca59f09b9470b5cfcb2c206fb3e (diff) | |
| download | gcc-038012e2d9211e56a2d6710a85314373af9d4d44.zip gcc-038012e2d9211e56a2d6710a85314373af9d4d44.tar.gz gcc-038012e2d9211e56a2d6710a85314373af9d4d44.tar.bz2 | |
[nvptx] Fix bar.sync position
2018-03-20 Tom de Vries <tom@codesourcery.com>
PR target/84952
* config/nvptx/nvptx.c (nvptx_single): Don't neuter bar.sync.
(nvptx_process_pars): Emit bar.sync asap and alap.
From-SVN: r258676
Diffstat (limited to 'gcc')
| -rw-r--r-- | gcc/ChangeLog | 6 | ||||
| -rw-r--r-- | gcc/config/nvptx/nvptx.c | 9 |
2 files changed, 12 insertions, 3 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 460802d..c1ac2f9 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,5 +1,11 @@ 2018-03-20 Tom de Vries <tom@codesourcery.com> + PR target/84952 + * config/nvptx/nvptx.c (nvptx_single): Don't neuter bar.sync. + (nvptx_process_pars): Emit bar.sync asap and alap. + +2018-03-20 Tom de Vries <tom@codesourcery.com> + PR target/84954 * config/nvptx/nvptx.c (prevent_branch_around_nothing): Also update seen_label if seen_label is already set. diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c index 7b0b182..1ba27e3 100644 --- a/gcc/config/nvptx/nvptx.c +++ b/gcc/config/nvptx/nvptx.c @@ -3969,7 +3969,9 @@ nvptx_single (unsigned mask, basic_block from, basic_block to) while (true) { /* Find first insn of from block. */ - while (head != BB_END (from) && !INSN_P (head)) + while (head != BB_END (from) + && (!INSN_P (head) + || recog_memoized (head) == CODE_FOR_nvptx_barsync)) head = NEXT_INSN (head); if (from == to) @@ -4018,6 +4020,7 @@ nvptx_single (unsigned mask, basic_block from, basic_block to) { default: break; + case CODE_FOR_nvptx_barsync: case CODE_FOR_nvptx_fork: case CODE_FOR_nvptx_forked: case CODE_FOR_nvptx_joining: @@ -4275,8 +4278,8 @@ nvptx_process_pars (parallel *par) nvptx_wpropagate (false, par->forked_block, par->forked_insn); nvptx_wpropagate (true, par->forked_block, par->fork_insn); /* Insert begin and end synchronizations. */ - emit_insn_after (nvptx_wsync (false), par->forked_insn); - emit_insn_before (nvptx_wsync (true), par->joining_insn); + emit_insn_before (nvptx_wsync (false), par->forked_insn); + emit_insn_before (nvptx_wsync (true), par->join_insn); } else if (par->mask & GOMP_DIM_MASK (GOMP_DIM_VECTOR)) nvptx_vpropagate (par->forked_block, par->forked_insn); |