OffTarget

Observed issues › stress/cce-c-flto-offload-dropped

Appears to behave incorrectlyCrashes or aborts at run time · C · OpenMP · Open

craycc -flto: the offload code is dropped; the program fails under OMP_TARGET_OFFLOAD=mandatory

Legality: contested. An independent adjudicating reading finds the result not required: The program is conforming. The specification says the initial thread may execute on the requested device if it is available and supported, and that with OMP_TARGET_OFFLOAD=mandatory a device that is unavailable or not supported by the implementation is treated as omp_invalid_device. The implementation (here with -flto) may not support offloading the region, so the abort is permitted. 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.0correctcorrectcorrectcorrectcorrectn/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/afails another wayfails another way
NVHPC 26.9n/an/an/an/an/acorrectcorrect
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 incorrectlyCCE 19.0.0, MI300A unified
CCE 19.0.0, MI300A copy
CCE 19.0.0, MI250X
show
ACC: OpenMP offload failed when OMP_TARGET_OFFLOAD=MANDATORY at address 0x200ac0
behaves incorrectlyCCE 20.0.2, MI300A unified
CCE 20.0.2, MI300A copy
CCE 20.0.2, MI250X
show
ACC: OpenMP offload failed when OMP_TARGET_OFFLOAD=MANDATORY at address 0x200a60
behaves incorrectlyCCE 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
show
ACC: OpenMP offload failed when OMP_TARGET_OFFLOAD=MANDATORY at address 0x200ab0
fails another wayNVHPC 25.11, A100 separate
NVHPC 25.11, A100 managed
show
nvc-Error-Unknown switch: -flto

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 -O2 -flto 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 -O2 -flto 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 -O2 -flto 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 -O2 -flto 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 -O2 -flto repro.c -o repro
HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro
Program (69 lines)
// # Finding: craycc -flto drops the OpenMP offload code
//
// - compilers: CCE craycc 19.0.0, 20.0.2, 21.0.3 (C), MI300A (gfx942) and MI250X (gfx90a)
// - build: the catalog's CCE C recipe (offload/engine/toolchains.py C_FLAGS) plus `-O2 -flto`;
//   without -flto the same file is correct on every CCE release; AFAR amdclang -flto is correct
// - expected: sum= 2016 dev= 1 (a target teams distribute loop fills a[i] = i on the device)
// - build flags: -O2 -flto
//   (read by the confirmer, offload.engine.confirm.programs.file_ref, which adds them to every
//   build of this reproducer; its -O0 control keeps -flto)
// - actual (OMP_TARGET_OFFLOAD=mandatory): the program builds, then exits 255 with "ACC: OpenMP
//   offload failed when OMP_TARGET_OFFLOAD=MANDATORY at address 0x200a...". Observed cells:
//   MI300A HSA_XNACK 1 (CCE 19/20/21, audit runs.jsonl), MI250X HSA_XNACK 0 (Tioga, CCE 19/20/21,
//   runs-tioga.jsonl), and MI300A HSA_XNACK 0 (CCE 19 and 21; CCE 19/20/21,
//   runs-fix2.jsonl). Under
//   OMP_TARGET_OFFLOAD=default the -flto binary runs the region on the host: "sum= 2016 dev= 0".
// - the offload code is absent (readelf/nm of the executables, CCE 19.0.0 and 21.0.3,
//   MI300A): without -flto the binary has a cray_omp_offloading_entries section, the
//   .offloading.entry.__omp_offloading_* symbols and 7 strings naming amdgcn/gfx942; with -O2 -flto
//   the entries section and the entry symbols are gone (only the .region_id symbols remain) and no
//   amdgcn/gfx942 string is left. AFAR amdclang -flto keeps .llvm.offloading and its entries and
//   runs on the device. Not run: the non--flto control on MI250X.
// - evidence: results/stress/audit/runs.jsonl and runs-tioga.jsonl (tools/campaigns/stress_audit.py on
//   data/reports/fuzz-stress-audit/r8-craycc-flto.c); the audit rerun's runs are committed as
//   data/reports/fuzz-stress-audit/rereview-1-lto.jsonl. In the stress sweeps every C program of
//   the `lto` variant failed the same way (s1-var, 3 CCE releases x 3 layouts), and every other
//   variant agreed
// - confirmer (which carries the "build flags" line above into every build): on CCE 19 in
//   both modes (`tools/confirm.py case <this file> --toolchain cce19 --mode unified|copy --expected
//   "sum= 2016 dev= 1" --run`, 2026-10-25) the verdict is needs_human (R22-expected-weak): CCE 19, 20
//   and 21 built with -O2 -flto exit 255 on every run and AFAR amdclang -O2 -flto prints the expected
//   value, but the expected value rests on the stated value alone (no sequential elision of this C
//   program), which the rules call provisional. Before the confirmer could not add -flto
//   (data/reports/fuzz-stress.md). The installed CCE documentation has no craycc page
//   and does not say that -flto is unsupported with offload; if HPE documents that, this is a
//   documented limitation, not a defect.
// - legality: legal C11 / OpenMP 5.0 by reading: a[0:64] is mapped from and each iteration writes
//   its own element; dev is mapped tofrom and written by the one initial thread of a target region
//   (amdclang -fopenmp -fsyntax-only, not of the family under test, accepts it)
/* Audit: craycc -flto with OpenMP offload. A target teams
   distribute loop; built with -flto it must still run on the device.
   Expected (with and without -flto): sum= 2016 dev= 1 */
#include <stdio.h>
#include <omp.h>
int main(void) {
  long long a[64], sum = 0;
  int dev = -1;
  #pragma omp target teams distribute map(from: a[0:64])
  for (int i = 0; i < 64; i++)
    a[i] = i;
  for (int i = 0; i < 64; i++)
    sum += a[i];
  #pragma omp target map(tofrom: dev)
  dev = omp_is_initial_device() ? 0 : 1;
  printf("sum= %lld dev= %d\n", sum, dev);
  return 0;
}
// spec (requirement): a target construct includes an executable region that is executed on a device, and with
//   OMP_TARGET_OFFLOAD=mandatory a construct that cannot use its device is an error instead of a host fallback. The -flto option
//   is a vendor build option that OpenMP does not mention, so there is no OpenMP text on link-time optimization itself; the
//   nearest normative text is the device execution of the target region.
// "The target construct provides a superset of the functionality provided by the target data directive, except for the
//   use_device_ptr and use_device_addr clauses. The functionality added to the target directive is the inclusion of an executable
//   region to be executed on a device." (5.2 §13.8 p.284 l.11-13)
// "The mandatory value specifies that the effect of any device construct or device memory routine that uses a device that is
//   unavailable or not supported by the implementation, or uses a non-conforming device number, is as if the omp_invalid_device
//   device number was used." (5.2 §21.2.8 p.610 l.28-30)
// spec (this program): the target teams distribute region must still execute on the device in a program built with -flto; the
//   binary has lost its offload entries, so OMP_TARGET_OFFLOAD=mandatory aborts (the default setting silently runs the region on
//   the host, dev= 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

sum= 2016 dev= 1

Observed where it behaves incorrectly

ACC: OpenMP offload failed when OMP_TARGET_OFFLOAD=MANDATORY at address 0x200ac0

What the specification requires

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

OpenMP 5.2 §13.8 p.284 l.11-13

The target construct provides a superset of the functionality provided by the target data directive, except for the use_device_ptr and use_device_addr clauses. The functionality added to the target directive is the inclusion of an executable region to be executed on a device.

OpenMP 5.2 §21.2.8 p.610 l.28-30

The mandatory value specifies that the effect of any device construct or device memory routine that uses a device that is unavailable or not supported by the implementation, or uses a non-conforming device number, is as if the omp_invalid_device device number was used.

Evidence

Same compiler, one change

  • built without -flto (-O2): AFAR 24.3.0: correct; CCE 19.0.0: correct; CCE 20.0.2: correct; CCE 21.0.3: correct; CCE 21.0.4: correct
    Control program
    #include <stdio.h>
    #include <omp.h>
    int main(void) {
      long long a[64], sum = 0;
      int dev = -1;
      #pragma omp target teams distribute map(from: a[0:64])
      for (int i = 0; i < 64; i++)
        a[i] = i;
      for (int i = 0; i < 64; i++)
        sum += a[i];
      #pragma omp target map(tofrom: dev)
      dev = omp_is_initial_device() ? 0 : 1;
      printf("sum= %lld dev= %d\n", sum, dev);
      return 0;
    }
    

Other compilers run it correctly

AFAR 24.3.0, NVHPC 26.9, GCC 16.2.0