[Bug target/104422] New: nvptx: for-3.exe fail with driver 390.x

[Bug target/104422] New: nvptx: for-3.exe fail with driver 390.x

[vries at gcc dot gnu.org]

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

            Bug ID: 104422
           Summary: nvptx: for-3.exe fail with driver 390.x
           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: ---

While testing libgomp using legacy driver 390.x on a maxwell card, Quadro K620
I ran into a for-3.exe execution failure.

I've minimized it to this:
...
$ cat for-3.c  
/* { dg-additional-options "-std=gnu99" { target c } } */

extern void abort ();

#pragma omp declare target

int a[1500];

void
test_ds_normal (void)
{
#pragma omp target update to(a)
  ;

#pragma omp target teams
#pragma distribute simd
  for (int i = 0; i < 1500; i++)
    a[i] += 2;

#pragma omp target update from(a)
    ;
}

#pragma omp end declare target

int
main (void)
{
  int err = 0;

  for (int i = 0; i < 1500; i++)
    a[i] = i - 25;

  test_ds_normal ();

  for (int i = 0; i < 1500; i++)
    if (a[i] != i - 23)
      err = 1;

 if (err)
    abort ();

  return 0;
}
...

This might be a driver problem, we know of problems that we reported to nvidia
that they're not going to fix in this driver.

Then again, it might not be a driver problem...

[Bug target/104422] nvptx: for-3.exe fail with driver 390.x

[roger at nextmovesoftware dot com]

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

Roger Sayle <roger at nextmovesoftware dot com> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
                 CC|                            |roger at nextmovesoftware dot com

--- Comment #1 from Roger Sayle <roger at nextmovesoftware dot com> ---
Is the problem reproduceable (with the legacy driver) when the omp #pragma are
removed, and the program compiled with -mainkernel and executed with
nvptx-none-run-single?  This eliminates the GOMP (interaction) as source of the
problem.

[Bug target/104422] nvptx: for-3.exe fail with driver 390.x

[vries at gcc dot gnu.org]

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

--- Comment #2 from Tom de Vries <vries at gcc dot gnu.org> ---
Hmm, I reran on a(In reply to Tom de Vries from comment #0)
> #pragma distribute simd

omp missing ... I need to reproduce this.

[Bug target/104422] nvptx: for-3.exe fail with driver 390.x

[vries at gcc dot gnu.org]

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

--- Comment #3 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #0)
> While testing libgomp using legacy driver 390.x on a maxwell card, Quadro
> K620 I ran into a for-3.exe execution failure.

Reproduced with 390.147 driver on pascal card, GT 1030, gcc build from master
commit 73f4a989b7f ("[nvptx] Fix 'main (int argc)' compilation").

Reproduces both with and without GOMP_NVPTX_JIT=-O0.

[Bug target/104422] nvptx: for-3.exe fail with driver 390.x

[vries at gcc dot gnu.org]

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

--- Comment #4 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #3)
> Reproduces both with and without GOMP_NVPTX_JIT=-O0.

Pff, that was an artefact of having bumped the default ptx isa to 6.3.

So, let's try again ...

Reproduced with 390.147 driver on pascal card, GT 1030, gcc build from master
commit "[nvptx] Unbreak build, add PTX_ISA_SM70".  Doens't reproduce with
GOMP_NVPTX_JIT=-O0.

Minimal test-case:
...
#define N 1500

#pragma omp declare target
int a[N];

__attribute__((noinline, noclone)) void
f3_ds_normal ()
{
  volatile int s3 = 1;
#pragma omp distribute simd 
  for (int i = 0; i < N; i += s3)
    a[i] += 7;
}
#pragma omp end declare target

int
main ()
{
  for (int i = 0; i < N; i++)
    a[i] = i - 29;

#pragma omp target update to(a)

#pragma omp target teams
  f3_ds_normal ();

#pragma omp target update from(a)

  for (int i = 0; i < N; i++)
    if (a[i] != i - 22)
      __builtin_abort ();

  return 0;
}
...

[Bug target/104422] nvptx: for-3.exe fail with driver 390.x

[vries at gcc dot gnu.org]

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

--- Comment #5 from Tom de Vries <vries at gcc dot gnu.org> ---
Still on GT1030, does not reproduce with 470.x, neither the minimal nor the
complete for-3.c.

[Bug target/104422] nvptx: for-3.exe fail with driver 390.x

[vries at gcc dot gnu.org]

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

--- Comment #6 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #5)
> Still on GT1030, does not reproduce with 470.x, neither the minimal nor the
> complete for-3.c.

And the same for 510.x.

So, I'm parking this for now.  This may well be the umul bug that nvidia
decided it didn't want to fix on the 390.x driver branch.

It may be something else, but we'll only known after reducing the test-case at
ptx level, which might be difficult and time-intensive for an openmp test-case.