[nvptx] Update default ptx isa to 6.3
authorTom de Vries <tdevries@suse.de>
Wed, 26 Jan 2022 13:17:40 +0000 (14:17 +0100)
committerTom de Vries <tdevries@suse.de>
Tue, 1 Feb 2022 18:28:52 +0000 (19:28 +0100)
commit8ff0669f6d1d6126b7c010da02fa6532abb5e1ca
tree18bf26420bad3519dcecb79ccff84814e9978e5b
parent57f971f99209cc950d7e706b7b52f4c9ef1d10b0
[nvptx] Update default ptx isa to 6.3

With the following example, minimized from parallel-dims.c:
...
int
main (void)
{
  int vectors_max = -1;
  #pragma acc parallel num_gangs (1) num_workers (1) copy (vectors_max)
  {
    for (int i = 0; i < 2; i++)
      for (int j = 0; j < 2; j++)
        #pragma acc loop vector reduction (max: vectors_max)
        for (int k = 0; k < 32; k++)
          vectors_max = k;
  }

  if (vectors_max != 31)
    __builtin_abort ();

  return 0;
}
...
I run into (T400, driver version 470.94):
...
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims.c \
  -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 \
  execution test
...
The FAIL does not happen with GOMP_NVPTX_JIT=-O0.

The problem seems to be that the shfl insns for the vector reduction are not
executed uniformly by the warp.  Enforcing this by using shfl.sync fixes the
problem.

Fix this by setting the ptx isa to 6.3 by default, which allows the use of
shfl.sync.

Tested on x86_64 with nvptx accelerator.

gcc/ChangeLog:

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

* config/nvptx/nvptx.opt (mptx): Set to PTX_VERSION_6_3 by default.
gcc/config/nvptx/nvptx.opt