craycc: dist_schedule(static) on target teams distribute parallel for crashes (SIGSEGV) at HSA_XNACK=1

Affected variants

"behaves incorrectly" / "not supported": this issue, confirmed in up to four environments (another build directory, another node, a repeat). "fails another way": a different failure, not counted as this issue. "unknown": not run there, or the runs disagree.

Messages

RUN_FAIL

Reproduce

CCE 19.0.0, MI300A unified

module load cce/19.0.0 craype-accel-amd-gfx942 rocm/6.4.3
module unload cray-mpich
craycc -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx942 -fopenmp-version=52 -O2 repro.c -o repro
HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro

CCE 20.0.2, MI300A unified

module load cce/20.0.2 craype-accel-amd-gfx942 rocm/6.4.3
module unload cray-mpich
craycc -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx942 -fopenmp-version=52 -O2 repro.c -o repro
HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro

CCE 21.0.3, MI300A unified

module load cce/21.0.3 craype-accel-amd-gfx942 rocm/7.2.0
module unload cray-mpich
craycc -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx942 -fopenmp-version=52 -O2 repro.c -o repro
HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro

CCE 21.0.4, MI300A unified

module load cce/21.0.4 craype-accel-amd-gfx942 rocm/7.2.0
module unload cray-mpich
craycc -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx942 -fopenmp-version=52 -O2 repro.c -o repro
HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro

AFAR 24.3.0, MI300A unified (runs correctly)

$AFAR/bin/amdclang -fopenmp --offload-arch=gfx942 -fopenmp-version=52 -fopenmp-version=52 -O2 repro.c -o repro
HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro

/* Finding (C): dist_schedule(static) without a chunk size on a
   combined target teams distribute parallel for crashes on CCE.

   Legal OpenMP 5.2 C (5.2 §11.6.1: dist_schedule(static) divides the
   iterations into chunks of about equal size, one per team).

   Expected: "64 of 64".
   Observed at HSA_XNACK=1, -O2: CCE craycc 19.0.0, 20.0.2 and 21.0.3:
   the program dies with SIGSEGV (also without OMP_TARGET_OFFLOAD=
   mandatory; CCE 19 also at -O0). At HSA_XNACK=0 (copy mode) CCE 19
   prints "64 of 64". Without the dist_schedule clause, or with
   dist_schedule(static, 16), it prints "64 of 64" at HSA_XNACK=1. AFAR amdclang:
   "64 of 64". CCE ftn runs the Fortran form (omp52-clause-dist_schedule)
   correctly.

   craycc -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa
          -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx942
          -fopenmp-version=52 -O2 */
#include <stdio.h>
int main(void) {
  int b[64];
  for (int i = 0; i < 64; i++) b[i] = 0;
  #pragma omp target teams distribute parallel for map(from: b) \
      num_teams(4) dist_schedule(static)
  for (int i = 0; i < 64; i++) b[i] = i + 1;
  int n = 0;
  for (int i = 0; i < 64; i++) n += b[i] == i + 1;
  printf("%d of 64\n", n);
  return 0;
}

Source in the repository

Something to add? Dispute its legality · Report it fixed in a newer release · Report a mistake on this page

What the specification requires

The specification text this program relies on and the compiler is expected to support.

Evidence