Observed issues › catalog/cce-c-recursive-function-ice
Appears to behave incorrectlyCompiler crash · C · OpenMP · Fixed in 21.0.3
craycc: compiler crash on a recursive function called in a target region
Legality: independently confirmed. 2 independent readings find the program legal and the result required. 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 | unknown not run | correct | correct | n/a | n/a |
| CCE 19.0.0 | behaves incorrectly | behaves incorrectly | unknown not run | unknown not run | unknown not run | n/a | n/a |
| CCE 20.0.2 | behaves incorrectly | behaves incorrectly | unknown not run | unknown not run | unknown not run | n/a | n/a |
| CCE 21.0.3 | correct | correct | unknown not run | unknown not run | unknown not run | n/a | n/a |
| CCE 21.0.4 | correct | correct | unknown not run | unknown not run | unknown not run | n/a | n/a |
| NVHPC 25.11 | n/a | n/a | n/a | n/a | n/a | correct | correct |
| 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 | showPLEASE submit a bug report to HPE at https://support.hpe.com/ and include the crash backtrace, preprocessed source, and associated run script. Stack dump: 0. Program arguments: /opt/cray/pe/cce/19.0.0/cce-clang/x86_64/bin/clang-19 -cc1 -triple amdgcn-amd-amdhsa -mllvm -cray-omp-opt-fork-call -mllvm -cray-omp-parallel-opt-call -mllvm -cray-openmp-rename- 1. <eof> parser at end of file 2. Optimizer 3. Running pass "cray-gpu-lower" on module "cce-c-recursive-function-ice.c" Stack dump without symbol names (ensure you have llvm-symbolizer in your PATH or set the environment var `LLVM_SYMBOLIZER_PATH` to point to it): 0 clang-19 0x0000000003f8b1c7 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) + 39 |
| behaves incorrectly | CCE 20.0.2, MI300A unified CCE 20.0.2, MI300A copy | showPLEASE submit a bug report to HPE at https://support.hpe.com/ and include the crash backtrace, preprocessed source, and associated run script. Stack dump: 0. Program arguments: /opt/cray/pe/cce/20.0.2/cce-clang/x86_64/bin/clang-20 -cc1 -triple amdgcn-amd-amdhsa -mllvm -cray-omp-opt-fork-call -mllvm -cray-omp-parallel-opt-call -mllvm -cray-openmp-rename- 1. <eof> parser at end of file 2. Optimizer 3. Running pass "cray-gpu-lower" on module "cce-c-recursive-function-ice.c" Stack dump without symbol names (ensure you have llvm-symbolizer in your PATH or set the environment var `LLVM_SYMBOLIZER_PATH` to point to it): 0 clang-20 0x0000000003ed4da7 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) + 39 |
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
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
Program (58 lines)
// # Draft finding: omp52-c-w-recursive-routine on cce19
//
// - compiler: cce19
// - toolchain: cce19; arch: gfx942; machine: tuolumne; memory mode: unified; opt: -O2
// - program: catalog omp52-c-w-recursive-routine (c, version 52, sha256 5032bc0dcb16068b)
// - expected: None (None; )
// - actual: ice (ice: clang: error: clang frontend command failed due to signal (use -v to see invocation))
// - verdict: compiler_bug by rule R10-compile-ice: an internal compiler error (or the compiler rejecting code it generated) on a legal program: always the compiler's defect
// - reason:
// - citations: RUL-COMPILE-TIME (rulings.md)
// - legality: legal: offload.engine.prerun.check accepts it (rerun: offload.engine.rules with the
// race rule and the restriction checks, and the AFAR clang front end, not under test here); no
// compiler under test, gfortran or gcc is a legality witness
// - witnesses: none
// - withdrawn: none
//
// Generated by compliance.widen (OpenMP 5.2, C11; catalog widening). Do not edit.
// id: omp52-c-w-recursive-routine
// primary atom: base.device.recursion
// atoms: base.device.recursion
// legal: yes
// observable: yes
// ran_on_host: inside the target region of interest
// evidence: construct
// needs: none
// std: c11
// feature: a recursive declare target function
// spec: OpenMP 5.2 §1.7, p.35: C11 is a normative base language
#include <stdio.h>
#include <stdbool.h>
#include <omp.h>
enum { n = 64 };
#pragma omp declare target
int sumto(int k) { return k <= 0 ? 0 : k + sumto(k - 1); }
#pragma omp end declare target
int main(void) {
int a[n], b[n], ref[n];
int host_flag = 1;
for (int i = 0; i < n; i++) {
a[i] = i + 1;
b[i] = 0;
ref[i] = ((i + 1) % 8) * ((i + 1) % 8 + 1) / 2;
}
#pragma omp target teams distribute parallel for map(to: a) map(from: b) map(tofrom: \
host_flag)
for (int i = 0; i < n; i++) {
b[i] = sumto(a[i] % 8);
if (i == 0) host_flag = omp_is_initial_device();
}
bool ok = true;
for (int i = 0; i < n; i++)
ok = ok && (b[i] == ref[i]);
printf("RESULT ran_on_host=%d check=%s\n", host_flag, ok ? "pass" : "fail");
return 0;
}
Something to add? Dispute its legality · Report it fixed in a newer release · Report a mistake on this page
Expected and observed
Observed where it behaves incorrectly
PLEASE submit a bug report to HPE at https://support.hpe.com/ and include the crash backtrace, preprocessed source, and associated run script. Stack dump: 0. Program arguments: /opt/cray/pe/cce/19.0.0/cce-clang/x86_64/bin/clang-19 -cc1 -triple amdgcn-amd-amdhsa -mllvm -cray-omp-opt-fork-call -mllvm -cray-omp-parallel-opt-call -mllvm -cray-openmp-rename- 1. <eof> parser at end of file 2. Optimizer 3. Running pass "cray-gpu-lower" on module "cce-c-recursive-function-ice.c" Stack dump without symbol names (ensure you have llvm-symbolizer in your PATH or set the environment var `LLVM_SYMBOLIZER_PATH` to point to it): 0 clang-19 0x0000000003f8b1c7 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) + 39
What the specification requires
The specification text this program relies on and the compiler is expected to support.
OpenMP 5.2 §1.7 p.35
Evidence
Same compiler, one change
- sumto computed without recursion (k*(k+1)/2): 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 <stdbool.h> #include <omp.h> enum { n = 64 }; #pragma omp declare target int sumto(int k) { return k * (k + 1) / 2; } #pragma omp end declare target int main(void) { int a[n], b[n], ref[n]; int host_flag = 1; for (int i = 0; i < n; i++) { a[i] = i + 1; b[i] = 0; ref[i] = ((i + 1) % 8) * ((i + 1) % 8 + 1) / 2; } #pragma omp target teams distribute parallel for map(to: a) map(from: b) map(tofrom: \ host_flag) for (int i = 0; i < n; i++) { b[i] = sumto(a[i] % 8); if (i == 0) host_flag = omp_is_initial_device(); } bool ok = true; for (int i = 0; i < n; i++) ok = ok && (b[i] == ref[i]); printf("RESULT ran_on_host=%d check=%s\n", host_flag, ok ? "pass" : "fail"); return 0; }
Other compilers run it correctly
AFAR 24.3.0, NVHPC 25.11, NVHPC 26.9, GCC 16.2.0