https://gcc.gnu.org/bugzilla/show_bug.cgi?id=85381

--- Comment #11 from Tom de Vries <vries at gcc dot gnu.org> ---
Author: vries
Date: Mon Jan  7 08:10:56 2019
New Revision: 267630

URL: https://gcc.gnu.org/viewcvs?rev=267630&root=gcc&view=rev
Log:
[nvptx] Don't emit barriers for empty loops -- fix

When compiling an empty loop:
...
  long long v1;
  #pragma acc parallel num_gangs (640) num_workers(1) vector_length (128)
  #pragma acc loop
    for (v1 = 0; v1 < 20; v1 += 2)
        ;
...
the compiler emits two subsequent bar.syncs.  This triggers some bug on my
quadro m1200 (I'm assuming in the ptxas/JIT compiler) that hangs the testcase.

This patch works around the bug by doing an optimization: we detect that this
is
an empty loop (a forked immediately followed by a joining), and don't emit the
barriers.

The patch does not include the test-case yet, since vector_length (128) is not
yet supported at this point.

2019-01-07  Tom de Vries  <tdevr...@suse.de>

        PR target/85381
        * config/nvptx/nvptx.c (nvptx_process_pars): Don't emit barriers for
        empty loops.

Modified:
    trunk/gcc/ChangeLog
    trunk/gcc/config/nvptx/nvptx.c

Reply via email to