Observed issues › stress/cce-c-flto-offload-dropped
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
| Compiler | MI300A 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.0 | correct | correct | correct | correct | correct | n/a | n/a |
| CCE 19.0.0 | behaves incorrectly | behaves incorrectly | behaves incorrectly | unknown not run | unknown not run | n/a | n/a |
| CCE 20.0.2 | behaves incorrectly | behaves incorrectly | behaves incorrectly | unknown not run | unknown not run | n/a | n/a |
| CCE 21.0.3 | behaves incorrectly | behaves incorrectly | behaves incorrectly | unknown not run | unknown not run | n/a | n/a |
| CCE 21.0.4 | behaves incorrectly | behaves incorrectly | behaves incorrectly | unknown not run | unknown not run | n/a | n/a |
| NVHPC 25.11 | n/a | n/a | n/a | n/a | n/a | fails another way | fails another way |
| NVHPC 26.9 | n/a | n/a | n/a | n/a | n/a | correct | correct |
| GCC 16.2.0 | correct | correct | correct | correct | correct | correct | n/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
| Result | Compiler and GPU | Message |
|---|---|---|
| behaves incorrectly | CCE 19.0.0, MI300A unified CCE 19.0.0, MI300A copy CCE 19.0.0, MI250X | showACC: OpenMP offload failed when OMP_TARGET_OFFLOAD=MANDATORY at address 0x200ac0 |
| behaves incorrectly | CCE 20.0.2, MI300A unified CCE 20.0.2, MI300A copy CCE 20.0.2, MI250X | showACC: OpenMP offload failed when OMP_TARGET_OFFLOAD=MANDATORY at address 0x200a60 |
| behaves incorrectly | 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 | showACC: OpenMP offload failed when OMP_TARGET_OFFLOAD=MANDATORY at address 0x200ab0 |
| fails another way | NVHPC 25.11, A100 separate NVHPC 25.11, A100 managed | shownvc-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).
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