Observed issues › nested/cce-c-struct-and-nested-pointee-sections-ice
CCE: compiler crash on nested pointee sections in one target construct
Legality: independently confirmed. 2 independent readings find the program legal and the result required. How this is checked
The CCE front end crashes on a target construct mapping nested structures through pointer members.
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 | behaves incorrectly | behaves incorrectly | unknown not run | unknown not run | unknown not run | n/a | n/a |
| CCE 21.0.4 | behaves incorrectly | behaves incorrectly | 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 | showclang-19: /workspace/llvm/include/llvm/Support/Casting.h:109: static bool llvm::isa_impl_cl<To, const From*>::doit(const From*) [with To = clang::FieldDecl; From = clang::ValueDecl]: Assertion `Val && 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 x86_64-unknown-linux-gnu -mllvm -cray-omp-opt-fork-call -mllvm -cray-omp-parallel-opt-call -mllvm -cray-openmp- 1. <eof> parser at end of file 2. Per-file LLVM IR generation 3. site-nested--cce-c-struct-and-nested-pointee-sections-ice.c:38:5: Generating code for declaration 'main' #0 0x0000000003f8b1c7 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/opt/cray/pe/cce/19.0.0/cce-clang/x86_64/bin/clang-19+0x3f8b1c7) |
| behaves incorrectly | CCE 20.0.2, MI300A unified CCE 20.0.2, MI300A copy | showclang-20: /workspace/llvm/include/llvm/Support/Casting.h:109: static bool llvm::isa_impl_cl<To, const From*>::doit(const From*) [with To = clang::FieldDecl; From = clang::ValueDecl]: Assertion `Val && 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/20.0.2/cce-clang/x86_64/bin/clang-20 -cc1 -triple x86_64-unknown-linux-gnu -mllvm -cray-omp-opt-fork-call -mllvm -cray-omp-parallel-opt-call -mllvm -cray-openmp- 1. <eof> parser at end of file 2. Per-file LLVM IR generation 3. site-nested--cce-c-struct-and-nested-pointee-sections-ice.c:38:5: Generating code for declaration 'main' #0 0x0000000003ed4da7 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/opt/cray/pe/cce/20.0.2/cce-clang/x86_64/bin/clang-20+0x3ed4da7) |
| behaves incorrectly | CCE 21.0.3, MI300A unified CCE 21.0.3, MI300A copy | showclang-21: /workspace/llvm/include/llvm/Support/Casting.h:109: static bool llvm::isa_impl_cl<To, const From*>::doit(const From*) [with To = clang::FieldDecl; From = clang::ValueDecl]: Assertion `Val && 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/21.0.3/cce-clang/x86_64/bin/clang-21 -cc1 -triple x86_64-unknown-linux-gnu -mllvm -cray-omp-opt-fork-call -mllvm -cray-omp-parallel-opt-call -mllvm -cray-openmp- 1. <eof> parser at end of file 2. Per-file LLVM IR generation 3. site-nested--cce-c-struct-and-nested-pointee-sections-ice.c:38:5: Generating code for declaration 'main' #0 0x0000000002bf69a7 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/opt/cray/pe/cce/21.0.3/cce-clang/x86_64/bin/clang-21+0x2bf69a7) |
| behaves incorrectly | CCE 21.0.4, MI300A unified CCE 21.0.4, MI300A copy | showclang-21: /workspace/llvm/include/llvm/Support/Casting.h:109: static bool llvm::isa_impl_cl<To, const From*>::doit(const From*) [with To = clang::FieldDecl; From = clang::ValueDecl]: Assertion `Val && 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/21.0.4/cce-clang/x86_64/bin/clang-21 -cc1 -triple x86_64-unknown-linux-gnu -mllvm -cray-omp-opt-fork-call -mllvm -cray-omp-parallel-opt-call -mllvm -cray-openmp- 1. <eof> parser at end of file 2. Per-file LLVM IR generation 3. site-nested--cce-c-struct-and-nested-pointee-sections-ice.c:38:5: Generating code for declaration 'main' #0 0x0000000002bf69b7 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/opt/cray/pe/cce/21.0.4/cce-clang/x86_64/bin/clang-21+0x2bf69b7) |
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 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 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 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 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 repro.c -o repro HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro
Program (66 lines)
// # Finding (nested-data): CCE craycc crashes in CodeGen: a target construct maps a structure and the structure and array
// # reached through an embedded structure and a pointer
//
// - toolchains and modes that fail: cce19 (CCE 19.0.0), cce20 (CCE 20.0.2), cce21 (CCE 21.0.3), cce21_4 (CCE 21.0.4); modes:
// unified, copy; -O2
// - build: craycc -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx942 -O2
// - failure: ice
// - observed: `PLEASE submit a bug report to HPE at https://support.hpe.com/ and include the crash backtrace, preprocessed source,
// and associated run script.`; identical signature on 8 of 8 failing cells
// - expected: compiles and prints `tot= 16` with `check=pass`
// - rule: "A structure and an element of it may both be list items" in the sense of 5.2 §5.8.3 p.156 l.6-7 ("unless one
// is the containing structure of the other"), and a section reached through pointer members is a list item of a map
// clause (§5.8.3 p.151 l.13); the construct maps r, then r.nx.nx[0:1] and r.nx.nx->a[0:8] through the embedded
// structure r.nx and its pointer member
// - same-compiler control: controls/j1-embedded-struct-pointer-to-struct.c (one change from this file; correct on cce19 (CCE
// 19.0.0), cce20 (CCE 20.0.2), cce21 (CCE 21.0.3), cce21_4 (CCE 21.0.4); rows
// results/findings-baseline/nested-j-controls.cce.jsonl, results/nested/red/tuo4.jsonl, results/nested/red/tuo5.jsonl,
// results/nested/red/wgs.jsonl, results/nested/red/wnv.jsonl); controls/j6-ctl-struct-and-one-pointee-level.c (one change from
// this file; correct on cce19 (CCE 19.0.0), cce20 (CCE 20.0.2), cce21 (CCE 21.0.3), cce21_4 (CCE 21.0.4); rows
// results/findings-baseline/nested-j-controls.cce.jsonl, results/nested/red/tuo5.jsonl, results/nested/red/wgs.jsonl,
// results/nested/red/wnv.jsonl); controls/j7-ctl-struct-and-embedded-pointee.c (one change from this file; correct on cce19 (CCE
// 19.0.0), cce20 (CCE 20.0.2), cce21 (CCE 21.0.3), cce21_4 (CCE 21.0.4); rows
// results/findings-baseline/nested-j-controls.cce.jsonl, results/nested/red/tuo5.jsonl, results/nested/red/wgs.jsonl,
// results/nested/red/wnv.jsonl); controls/j2-ctl-pointer-instead-of-embedded.c (one change from this file; correct on cce19 (CCE
// 19.0.0), cce20 (CCE 20.0.2), cce21 (CCE 21.0.3), cce21_4 (CCE 21.0.4); rows
// results/findings-baseline/nested-j-controls.cce.jsonl, results/nested/red/tuo5.jsonl, results/nested/red/wgs.jsonl,
// results/nested/red/wnv.jsonl)
// - other families: afar (AFAR 24.3.0), nvhpc (NVHPC 25.11), nvhpc269 (NVHPC 26.9), gcc_g90 (GCC 16.2.0), gcc_sm80 (GCC 16.2.0) run
// the file correctly (rows results/nested/red/tuo5.jsonl, results/nested/red/wgs.jsonl, results/nested/red/wnv.jsonl, site rows,
// check case site-nested--cce-c-struct-and-nested-pointee-sections-ice)
// - gate and confirmer: campaign gate of the programs below: 40 cells, 40 confirmed in the four environments (a-d; (c) on another
// node or GPU; rows results/nested/big/gate_reruns.part.<host>.3222831.jsonl,
// results/nested/big/gate_reruns.part.<host>.1040717.jsonl, results/nested/big/gate_reruns.part.<host>.1284552.jsonl,
// results/nested/big/gate_reruns.part.<host>.1251749.jsonl, results/nested/big/gate_reruns.part.<host>.1517583.jsonl,
// results/nested/big/gate_reruns.part.<host>.1434668.jsonl, results/nested/big/rows.part.<host>.1956389.jsonl,
// results/nested/big/rows.part.<host>.1523848.jsonl, results/nested/big/rows.part.<host>.1239313.jsonl,
// results/nested/big/rows.part.<host>.1197409.jsonl, results/nested/big/rows.part.<host>.1146311.jsonl,
// results/nested/main/gate_reruns.part.<host>.961942.jsonl, results/nested/main/gate_reruns.part.<host>.919396.jsonl,
// results/nested/main/gate_reruns.part.<host>.3910568.jsonl,
// results/nested/main/gate_reruns.part.<host>.4072524.jsonl, results/nested/main/rows.part.<host>.994167.jsonl)
// - programs (5): nd-c-omp-always-chain-d2-53712162, nd-c-omp-always-chain-d3-71393589, nd-c-omp-close-chain-d3-87bce034,
// nd-c-omp-one-chain-d3-9f7ac271, nd-c-omp-one-chain-d3-d6c766c9
// - legality: the rules (offload.engine.rules) accept it; AFAR's clang compiles and runs it (a front end of another family)
#include <stdio.h>
#include <stdlib.h>
#include <omp.h>
struct n2 { int *a; };
struct n1 { struct n2 *nx; };
struct n0 { int v; struct n1 nx; };
int main(void)
{
int hf = 1, i, tot = 0, ok;
struct n0 r;
r.v = 0; r.nx.nx = calloc(1, sizeof *r.nx.nx); r.nx.nx->a = malloc(8 * sizeof(int));
for (i = 0; i < 8; i++) r.nx.nx->a[i] = 2;
#pragma omp target map(from: hf)
hf = omp_is_initial_device() ? 1 : 0;
#pragma omp target map(tofrom: tot) map(tofrom: r) map(tofrom: r.nx.nx[0:1]) map(tofrom: r.nx.nx->a[0:8])
{
for (int q = 0; q < 8; q++) tot += r.nx.nx->a[q];
}
ok = tot == 16;
printf("tot= %d\n", tot);
printf("RESULT ran_on_host=%d check=%s\n", hf, 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
clang-19: /workspace/llvm/include/llvm/Support/Casting.h:109: static bool llvm::isa_impl_cl<To, const From*>::doit(const From*) [with To = clang::FieldDecl; From = clang::ValueDecl]: Assertion `Val && 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 x86_64-unknown-linux-gnu -mllvm -cray-omp-opt-fork-call -mllvm -cray-omp-parallel-opt-call -mllvm -cray-openmp- 1. <eof> parser at end of file 2. Per-file LLVM IR generation 3. site-nested--cce-c-struct-and-nested-pointee-sections-ice.c:38:5: Generating code for declaration 'main' #0 0x0000000003f8b1c7 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/opt/cray/pe/cce/19.0.0/cce-clang/x86_64/bin/clang-19+0x3f8b1c7)
What the specification requires
The specification text this program relies on and the compiler is expected to support.
OpenMP 5.2 §5.8.3 p.156 l.6-7
OpenMP 5.2 §5.8.3 p.151 l.13
Evidence
Same compiler, one change
- one change from the finding; correct on cce19, cce20, cce21, cce21_4: not run yet
Control program
// a structure member that is a structure holding a pointer to a structure that holds a pointer: map(tofrom: r.nx.nx[0:1]) map(tofrom: r.nx.nx->a[0:8]) #include <stdio.h> #include <stdlib.h> #include <omp.h> struct n2 { int *a; }; struct n1 { struct n2 *nx; }; struct n0 { int v; struct n1 nx; }; int main(void) { int hf = 1, i, tot = 0, ok; struct n0 r; r.v = 0; r.nx.nx = calloc(1, sizeof *r.nx.nx); r.nx.nx->a = malloc(8 * sizeof(int)); for (i = 0; i < 8; i++) r.nx.nx->a[i] = 2; #pragma omp target map(from: hf) hf = omp_is_initial_device() ? 1 : 0; #pragma omp target map(tofrom: tot) map(tofrom: r.nx.nx[0:1]) map(tofrom: r.nx.nx->a[0:8]) { for (int q = 0; q < 8; q++) tot += r.nx.nx->a[q]; } ok = tot == 16; printf("tot= %d\n", tot); printf("RESULT ran_on_host=%d check=%s\n", hf, ok ? "pass" : "fail"); return 0; } - one change from the finding; correct on cce19, cce20, cce21, cce21_4: not run yet
Control program
// a structure member that is a structure holding a pointer to a structure that holds a pointer: map(tofrom: r) map(tofrom: r.nx[0:1]) map(tofrom: r.nx->a[0:8]) #include <stdio.h> #include <stdlib.h> #include <omp.h> struct n2 { int *a; }; struct n0 { int v; struct n2 *nx; }; int main(void) { int hf = 1, i, tot = 0, ok; struct n0 r; r.v = 0; r.nx = calloc(1, sizeof *r.nx); r.nx->a = malloc(8 * sizeof(int)); for (i = 0; i < 8; i++) r.nx->a[i] = 2; #pragma omp target map(from: hf) hf = omp_is_initial_device() ? 1 : 0; #pragma omp target map(tofrom: tot) map(tofrom: r) map(tofrom: r.nx[0:1]) map(tofrom: r.nx->a[0:8]) { for (int q = 0; q < 8; q++) tot += r.nx->a[q]; } ok = tot == 16; printf("tot= %d\n", tot); printf("RESULT ran_on_host=%d check=%s\n", hf, ok ? "pass" : "fail"); return 0; } - one change from the finding; correct on cce19, cce20, cce21, cce21_4: not run yet
Control program
// a structure member that is a structure holding a pointer to a structure that holds a pointer: map(tofrom: r) map(tofrom: r.nx.s.a[0:8]) #include <stdio.h> #include <stdlib.h> #include <omp.h> struct n2 { int *a; }; struct n1 { struct n2 s; }; struct n0 { int v; struct n1 nx; }; int main(void) { int hf = 1, i, tot = 0, ok; struct n0 r; r.v = 0; r.nx.s.a = malloc(8 * sizeof(int)); for (i = 0; i < 8; i++) r.nx.s.a[i] = 2; #pragma omp target map(from: hf) hf = omp_is_initial_device() ? 1 : 0; #pragma omp target map(tofrom: tot) map(tofrom: r) map(tofrom: r.nx.s.a[0:8]) { for (int q = 0; q < 8; q++) tot += r.nx.s.a[q]; } ok = tot == 16; printf("tot= %d\n", tot); printf("RESULT ran_on_host=%d check=%s\n", hf, ok ? "pass" : "fail"); return 0; } - one change from the finding; correct on cce19, cce20, cce21, cce21_4: not run yet
Control program
// control: the middle structure reached through a pointer #include <stdio.h> #include <stdlib.h> #include <omp.h> struct n2 { int *a; }; struct n1 { struct n2 *nx; }; struct n0 { int v; struct n1 *nx; }; int main(void) { int hf = 1, i, tot = 0, ok; struct n0 r; r.v = 0; r.nx = calloc(1, sizeof *r.nx); r.nx->nx = calloc(1, sizeof *r.nx->nx); r.nx->nx->a = malloc(8 * sizeof(int)); for (i = 0; i < 8; i++) r.nx->nx->a[i] = 2; #pragma omp target map(from: hf) hf = omp_is_initial_device() ? 1 : 0; #pragma omp target map(tofrom: tot) map(tofrom: r.nx[0:1]) map(tofrom: r.nx->nx[0:1]) map(tofrom: r.nx->nx->a[0:8]) { for (int q = 0; q < 8; q++) tot += r.nx->nx->a[q]; } ok = tot == 16; printf("tot= %d\n", tot); printf("RESULT ran_on_host=%d check=%s\n", hf, ok ? "pass" : "fail"); return 0; }
Other compilers run it correctly
AFAR 24.3.0, NVHPC 25.11, NVHPC 26.9, GCC 16.2.0