OffTarget

Observed issues › stress/afar-c-lastprivate-combined-target-not-mapped

Appears to behave incorrectlyWrong answer · C · OpenMP · Open

Lastprivate on target teams distribute parallel for does not imply map(tofrom), value lost

Legality: independently confirmed. 2 independent readings find the program legal and the result required. How this is checked

Affected variants

CompilerMI300A gfx942
unified memory (XNACK 1)
MI300A gfx942
copy (XNACK 0)
MI250X gfx90a
copy (XNACK 0)
MI300X gfx942
copy (XNACK 0)
MI350X gfx950
copy (XNACK 0)
A100 sm_80
separate memory
A100 sm_80
managed memory
AFAR 24.3.0behaves incorrectlybehaves incorrectlybehaves incorrectlybehaves incorrectlybehaves incorrectlyn/an/a
CCE 19.0.0behaves incorrectlybehaves incorrectlybehaves incorrectlyunknown
not run
unknown
not run
n/an/a
CCE 20.0.2behaves incorrectlybehaves incorrectlybehaves incorrectlyunknown
not run
unknown
not run
n/an/a
CCE 21.0.3behaves incorrectlybehaves incorrectlybehaves incorrectlyunknown
not run
unknown
not run
n/an/a
CCE 21.0.4behaves incorrectlybehaves incorrectlybehaves incorrectlyunknown
not run
unknown
not run
n/an/a
NVHPC 25.11n/an/an/an/an/abehaves incorrectlybehaves incorrectly
NVHPC 26.9n/an/an/an/an/abehaves incorrectlybehaves incorrectly
GCC 16.2.0correctcorrectcorrectcorrectcorrectcorrectn/a

"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

ResultCompiler and GPUMessage
behaves incorrectlyAFAR 24.3.0, MI300A unified
AFAR 24.3.0, MI300A copy
AFAR 24.3.0, MI250X
AFAR 24.3.0, MI300X
AFAR 24.3.0, MI350X
CCE 19.0.0, MI300A unified
CCE 19.0.0, MI300A copy
CCE 19.0.0, MI250X
CCE 20.0.2, MI300A unified
CCE 20.0.2, MI300A copy
CCE 20.0.2, MI250X
CCE 21.0.3, MI300A unified
CCE 21.0.3, MI300A copy
CCE 21.0.3, MI250X
CCE 21.0.4, MI300A unified
CCE 21.0.4, MI300A copy
CCE 21.0.4, MI250X
NVHPC 25.11, A100 separate
NVHPC 25.11, A100 managed
NVHPC 26.9, A100 separate
NVHPC 26.9, A100 managed
show
l= 0 dev= 1

Reproduce

AFAR 24.3.0, MI300A unified

$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

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

NVHPC 25.11, A100 separate

export LD_LIBRARY_PATH=$NVHPC_25_11/compilers/lib
nvc -mp=gpu -gpu=cc80 -gpu=mem:separate -O2 repro.c -o repro
OMP_TARGET_OFFLOAD=mandatory ./repro

NVHPC 26.9, A100 separate

export LD_LIBRARY_PATH=$NVHPC_26_9/compilers/lib
nvc -mp=gpu -gpu=cc80 -gpu=mem:separate -O2 repro.c -o repro
OMP_TARGET_OFFLOAD=mandatory ./repro

GCC 16.2.0, MI300A unified (runs correctly)

export LD_LIBRARY_PATH=$GCC_16_2/lib64:/opt/rocm-7.2.1/lib
$GCC_16_2/bin/gcc -fopenmp -foffload=amdgcn-amdhsa -foffload-options=amdgcn-amdhsa=-march=gfx942 -O2 repro.c -o repro
HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro
Program (68 lines)
/* # Finding: AFAR amdclang drops lastprivate's implicit map on a combined
   # target construct (C)

   - compiler: AFAR amdclang 24.0.0git (drop 24.3.0), -fopenmp --offload-arch=gfx942 (MI300A, HSA_XNACK
     1 and 0) and gfx90a (MI250X, Tioga, XNACK 0), -O2
   - expected: l= 2560 dev= 1. lastprivate gives l the value of the sequentially last iteration
     (i = 64: 64 * 40) (5.2 §5.4.5), and "If the target construct is among the constituent constructs
     and the list item is not the base variable ... of a list item that appears in a map clause, the
     effect of the lastprivate clause is as if the same list item appears in a map clause with a
     map-type of tofrom" (5.2 §17.2 p.340 l.27-30; 5.0 §2.14 p.219 l.18-20)
   - actual: l= 0 dev= 1 (the region ran on the device; the final value was not copied back), on
     MI300A at XNACK 1 and 0 and on MI250X. The same with target teams distribute lastprivate(l)
     (lp-c-tdist.c).
   - XNACK 1 hang check: the audit saw this program and the other unmapped forms time out
     (rc 124 at a 60 s run limit) at XNACK 1 on <host>. It was rerun with a 600 s limit
     (tools/jobs/run_stress_fix2.sh b; results/stress/audit/runs-fix2-hang.jsonl) on <host>, 1031
     and 1011 (this program and lp-c-tdist.c), and the part-a run on 1014 had a 120 s limit. Fix
     Reran the other two forms that had timed out, target teams distribute simd and target
     parallel for (rereview-1-src/lp-tdsimd.c, lp-tpfor.c), with a 600 s limit on <host> and 1015
     (run_stress_fix2.sh d; runs-fix3-hang.jsonl). Every run ended in seconds with l= 0 at XNACK 1
     and 0. The timeout did not reproduce, so the observed outcome is l= 0; the audit's
     timeouts are recorded as unexplained.
   - same-compiler control: the same program with map(tofrom: l)
     (data/reports/fuzz-stress-audit/r2-lastprivate-combined-control.c): 2560 on MI300A (XNACK 1 and 0,
     every node above) and MI250X. Also 2560: target teams distribute lastprivate(l) map(tofrom: l)
     and target parallel for lastprivate(l) map(tofrom: l).
     : the control passes because l is long long. With an int item the
     explicit-map form fails too (0 or a garbage value), a second AFAR defect of items smaller than
     8 bytes (afar-c-lastprivate-distribute-parallel-for-small-type.c); this finding's own form
     loses the value with int and long long alike (lp= -1 from -1, both GPUs,
     data/evidence/lead_probe/leads/lp_int_nomap.c, lp_ll_nomap.c), so it stands as stated for its long long
     program and control.
   - no other family can witness:
       CCE craycc 19/20/21 give l= 0 on this program and also on the explicit-map control
       (stress/cce-c-distribute-lastprivate-lost.c, cce-c-parallel-for-in-target.c).
       NVHPC: NVIDIA's guide documents lastprivate on distribute as unsupported (HPC Compilers
       User's Guide 25.11, §7.4, 2.9.4 "The lastprivate clause is not supported"; 2.14 "Clauses on
       combined constructs are supported to the extent that the clauses are supported on the
       component constructs"), so nvc is not expected to implement this program; it gives l= 0
       with and without the map (stress/nvhpc-c-lastprivate-teams-distribute-parallel-for-lost.c).
       Its target teams distribute lastprivate(l) gives 0 without a map and 2560 with one, but that
       clause is documented as unsupported, so it is no evidence about the C rule.
       Fix (the optional check): on target parallel for, where lastprivate is documented
       (2.9.2), nvc 25.11 and 26.9 give l= 0 without a map (rereview-1-src/lp-tpfor.c,
       runs-nvhpc-fix3.jsonl, both memory modes) and 2560 with one: nvc misses the implicit map
       there too, so it is not a witness.
     So no C front end gives the expected value for an unmapped lastprivate on a combined target
     construct. The expected value rests on the spec text, AFAR's own explicit-map control, and the
     same rule in Fortran, which AFAR flang and CCE ftn implement (r2f-lastprivate-combined.f90:
     2560).
   - legality: legal by reading: the for loop variable is private; only the lastprivate item l is
     written in the loop; dev is written by the initial thread of the second target region
   - evidence: results/stress/audit/runs.jsonl, runs-tioga.jsonl, runs-fix2.jsonl,
     runs-fix2-hang.jsonl; audit rerun: data/reports/fuzz-stress-audit/rereview-1-lp.jsonl. The
     confirmer gave needs_human (R22: the CCE cells are wrong too, so it has no witness) */
#include <stdio.h>
#include <omp.h>
int main(void) {
  long long l = 0;
  int dev = -1;
  #pragma omp target teams distribute parallel for lastprivate(l)
  for (int i = 1; i <= 64; i++)
    l = (long long) i * 40;
  #pragma omp target map(tofrom: dev)
  dev = omp_is_initial_device() ? 0 : 1;
  printf("l= %lld dev= %d\n", l, dev);
  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

Expected and observed

Expected

l= 2560 dev= 1

Observed where it behaves incorrectly

l= 0 dev= 1

What the specification requires

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

OpenMP 5.2 §5.4.5

OpenMP 5.2 §17.2 p.340 l.27-30

If the target construct is among the constituent constructs and the list item is not the base variable ... of a list item that appears in a map clause, the effect of the lastprivate clause is as if the same list item appears in a map clause with a map-type of tofrom

OpenMP 5.0 §2.14 p.219 l.18-20

Evidence

Same compiler, one change

  • map(tofrom: l) added: AFAR 24.3.0: correct; CCE 19.0.0: behaves incorrectly; CCE 20.0.2: behaves incorrectly; CCE 21.0.3: behaves incorrectly; CCE 21.0.4: behaves incorrectly
    Control program
    #include <stdio.h>
    #include <omp.h>
    int main(void) {
      long long l = 0;
      int dev = -1;
      #pragma omp target teams distribute parallel for lastprivate(l) map(tofrom: l)
      for (int i = 1; i <= 64; i++)
        l = (long long) i * 40;
      #pragma omp target map(tofrom: dev)
      dev = omp_is_initial_device() ? 0 : 1;
      printf("l= %lld dev= %d\n", l, dev);
      return 0;
    }
    

Other compilers run it correctly

GCC 16.2.0