OffTarget

Observed issues › nested/cce-c-struct-and-nested-pointee-sections-ice

Appears to behave incorrectlyCompiler crash · C · OpenMP · Open

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

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.0correctcorrectunknown
not run
correctcorrectn/an/a
CCE 19.0.0behaves incorrectlybehaves incorrectlyunknown
not run
unknown
not run
unknown
not run
n/an/a
CCE 20.0.2behaves incorrectlybehaves incorrectlyunknown
not run
unknown
not run
unknown
not run
n/an/a
CCE 21.0.3behaves incorrectlybehaves incorrectlyunknown
not run
unknown
not run
unknown
not run
n/an/a
CCE 21.0.4behaves incorrectlybehaves incorrectlyunknown
not run
unknown
not run
unknown
not run
n/an/a
NVHPC 25.11n/an/an/an/an/acorrectcorrect
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
show
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)
behaves incorrectlyCCE 20.0.2, MI300A unified
CCE 20.0.2, MI300A copy
show
clang-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 incorrectlyCCE 21.0.3, MI300A unified
CCE 21.0.3, MI300A copy
show
clang-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 incorrectlyCCE 21.0.4, MI300A unified
CCE 21.0.4, MI300A copy
show
clang-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;
}

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

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