[Bug target/100428] New: [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test

[Bug target/100428] New: [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test

[vries at gcc dot gnu.org]

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

            Bug ID: 100428
           Summary: [nvptx, GOMP_NVPTX_JIT=-O0] FAIL:
                    libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-
                    7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0
                    -foffload=nvptx-none  -O0  execution test
           Product: gcc
           Version: 12.0
            Status: UNCONFIRMED
          Severity: normal
          Priority: P3
         Component: target
          Assignee: unassigned at gcc dot gnu.org
          Reporter: vries at gcc dot gnu.org
  Target Milestone: ---

When running libgomp testsuite using GOMP_NVPTX_JIT=-O0, I run into:
...
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none  -O0 
execution test
...

In more detail:
...
Execution timeout is: 300
spawn [open ...]^M
reduction-7.exe:
/home/vries/oacc/trunk/source-gcc/libgomp/testsuite/libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c:312:
v_p_2: Assertion `out[j * 32 + i] == (i + j) * 2' failed.
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none  -O0 
execution test
...

Using driver 450.119.03.

[Bug target/100428] [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test

[vries at gcc dot gnu.org]

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

--- Comment #1 from Tom de Vries <vries at gcc dot gnu.org> ---
Likewise, c++:
...
FAIL: libgomp.oacc-c++/../libgomp.oacc-c-c++-common/reduction-7.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none  -O0 
execution test
...

[Bug target/100428] [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test

[vries at gcc dot gnu.org]

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

Tom de Vries <vries at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
             Target|                            |nvptx

--- Comment #2 from Tom de Vries <vries at gcc dot gnu.org> ---
Minimized to:
...
#include <stdio.h>
#include <assert.h>

int
main (void)
{
  unsigned int i, j;
  unsigned int out;

  #pragma acc parallel num_gangs(1) num_workers(1) vector_length(32) \
                       copyout(out)
  {
    for (j = 1; j < 2; j++)
      {
        unsigned int acc = 0;

        #pragma acc loop vector reduction(+:acc)
        for (i = 0; i < 32; i++)
          acc += 1;

        out = acc;
      }
  }

  printf ("out: %u\n", out);
  fflush (0);
  assert (out == 32);

  return 0;
}
...
which gives us:
...
Execution timeout is: 300
spawn [open ...]^M
out: 1
reduction-7.exe: reduction-7.c:27: main: Assertion `out == 32' failed.
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none  -O0 
execution test
...

[Bug target/100428] [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test

[vries at gcc dot gnu.org]

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

--- Comment #3 from Tom de Vries <vries at gcc dot gnu.org> ---
The ptx code looks a lot like the cuda reproducer in PR99932 comment 4, so I'm
going to retest this once I get a driver where that one is fixed.

[Bug target/100428] [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test

[vries at gcc dot gnu.org]

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

--- Comment #4 from Tom de Vries <vries at gcc dot gnu.org> ---
FTR, reproduces with driver version 470.86 on Quadro M1200 and GeForce GT 1030.

[Bug target/100428] [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test

[cvs-commit at gcc dot gnu.org]

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

--- Comment #5 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Tom de Vries <vries@gcc.gnu.org>:

https://gcc.gnu.org/g:456de10c549379b74d4858f00d4b8817035a73fc

commit r12-6967-g456de10c549379b74d4858f00d4b8817035a73fc
Author: Tom de Vries <tdevries@suse.de>
Date:   Sun Jan 23 06:42:24 2022 +0100

    [nvptx] Handle nop in prevent_branch_around_nothing

    When running libgomp test-case reduction-7.c on an nvptx accelerator
    (T400, driver version 470.86) and GOMP_NVPTX_JIT=-O0, I run into:
    ...
    reduction-7.exe:reduction-7.c:312: v_p_2: \
      Assertion `out[j * 32 + i] == (i + j) * 2' failed.
    FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c \
      -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none \
      -O0  execution test
    ...

    During investigation I found ptx code like this:
    ...
    @ %r163 bra $L262;
    $L262:
    ...

    There's a known problem with executing this type of code, and a workaround
is
    in place to address this: prevent_branch_around_nothing.  The workaround
does
    not trigger though because it doesn't handle the nop insn.

    Fix this by handling the nop insn in prevent_branch_around_nothing.

    Tested libgomp on x86_64 with nvptx accelerator.

    gcc/ChangeLog:

    2022-01-27  Tom de Vries  <tdevries@suse.de>

            PR target/100428
            * config/nvptx/nvptx.cc (prevent_branch_around_nothing): Handle nop
            insn.

[Bug target/100428] [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test

[vries at gcc dot gnu.org]

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

Tom de Vries <vries at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
         Resolution|---                         |FIXED
             Status|UNCONFIRMED                 |RESOLVED
   Target Milestone|---                         |12.0

--- Comment #6 from Tom de Vries <vries at gcc dot gnu.org> ---
Fixed in gcc-12.