Observed issues › catalog/cce-c-map-iterator-ice
OpenMP 5.1/5.2/6.0 forms HPE does not claim are rejected, and some crash the compiler instead of a diagnostic
Legality: contested. An independent adjudicating reading finds the program not legal: Both readings are right. With k=1:2 the section w[k:n/2-(k+1)+1] gives w[1:31] and w[2:30], two instances with the same containing array w. The 5.2 map restriction forbids this, so the program is non-conforming and a compiler crash or rejection 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 | fails another way | fails another way | fails another way | fails another way | 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 | fails another way | fails another way |
| 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 | showDeclRefExpr for Decl not entered in LocalDeclMap? UNREACHABLE executed at /workspace/clang/lib/CodeGen/CGExpr.cpp:3180! 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. cce-c-map-iterator-ice.c:7:5: Generating code for declaration 'main' |
| behaves incorrectly | CCE 20.0.2, MI300A unified CCE 20.0.2, MI300A copy CCE 20.0.2, MI250X | showDeclRefExpr for Decl not entered in LocalDeclMap? UNREACHABLE executed at /workspace/clang/lib/CodeGen/CGExpr.cpp:3224! 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. cce-c-map-iterator-ice.c:7:5: Generating code for declaration 'main' |
| behaves incorrectly | CCE 21.0.3, MI300A unified CCE 21.0.3, MI300A copy CCE 21.0.3, MI250X | showDeclRefExpr for Decl not entered in LocalDeclMap? UNREACHABLE executed at /workspace/clang/lib/CodeGen/CGExpr.cpp:3331! 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. cce-c-map-iterator-ice.c:7:5: Generating code for declaration 'main' |
| behaves incorrectly | CCE 21.0.4, MI300A unified CCE 21.0.4, MI300A copy CCE 21.0.4, MI250X | showDeclRefExpr for Decl not entered in LocalDeclMap? UNREACHABLE executed at /workspace/clang/lib/CodeGen/CGExpr.cpp:3331! 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. cce-c-map-iterator-ice.c:7:5: Generating code for declaration 'main' |
| fails another way | AFAR 24.3.0, MI300A copy AFAR 24.3.0, MI250X AFAR 24.3.0, MI300X AFAR 24.3.0, MI350X | showRESULT ran_on_host=0 check=fail |
| fails another way | NVHPC 25.11, A100 separate NVHPC 25.11, A100 managed NVHPC 26.9, A100 separate NVHPC 26.9, A100 managed | show"cce-c-map-iterator-ice.c", line 26: error: invalid text in pragma
#pragma omp target enter data map(iterator(k = 1:2), to: w[k:n / 2 - (k + 1) + 1])
^
"cce-c-map-iterator-ice.c", line 26: error: invalid text in pragma
#pragma omp target enter data map(iterator(k = 1:2), to: w[k:n / 2 - (k + 1) + 1])
^
"cce-c-map-iterator-ice.c", line 26: error: extra text after expected end of preprocessing directive
#pragma omp target enter data map(iterator(k = 1:2), to: w[k:n / 2 - (k + 1) + 1]) |
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
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 -fopenmp-version=52 -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 -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 (86 lines)
// # Draft finding: omp52-c-clause-map-iterator on cce19
//
// - compiler: cce19
// - toolchain: cce19; arch: gfx942; machine: tuolumne; memory mode: unified; opt: -O2
// - program: catalog omp52-c-clause-map-iterator (c, version 52, sha256 a7be9001f201c65b)
// - expected: None (None; )
// - actual: ice (ice: PLEASE submit a bug report to HPE at https://support.hpe.com/ and include the crash backtrace, preprocessed source, and associated run script.)
// - 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 (legal by construction: offload.engine.gen / c_gen / acc_gen check every program against the catalog rules and the offload scope as it is built (no compiler))
// - witnesses: none
// - withdrawn: none
//
// Generated by compliance.c_gen (OpenMP 5.2, C11): the C port (offload.engine.f2c) of omp52-clause-map-iterator. Do not edit.
// id: omp52-c-clause-map-iterator
// primary atom: omp.clause.map.iterator
// legal: yes
// observable: yes
// ran_on_host: inside the target region of interest
// evidence: copy_mode_only
// needs: none
// std: c11
#include <stdio.h>
#include <stdbool.h>
#include <omp.h>
enum { n = 64 };
int main(void) {
int a[n];
int b[n];
int ref[n];
int host_flag;
bool ok;
int k;
int w[n];
bool pres;
bool shared_mem;
int zc[1];
for (int i = 1; i <= n; i++) {
a[i - 1] = i;
b[i - 1] = 0;
ref[i - 1] = 2 * i + 1;
}
host_flag = 1;
for (int k1_ = 0; k1_ < n; k1_++)
w[k1_] = 0;
#pragma omp target enter data map(iterator(k = 1:2), to: w[k:n / 2 - (k + 1) + 1])
pres = true;
for (int k2_ = 0; k2_ < 1; k2_++)
zc[k2_] = 0;
#pragma omp target map(to: zc)
{
zc[0] = 1;
}
shared_mem = zc[0] == 1;
#pragma omp target map(to: a) map(from: b) map(tofrom: host_flag)
{
host_flag = (omp_is_initial_device() ? 1 : 0);
for (int i = 1; i <= n; i++) {
b[i - 1] = 2 * a[i - 1] + 1;
}
}
if (!shared_mem) pres = pres && (omp_target_is_present(&w[1], omp_get_default_device()) != 0);
if (!shared_mem) pres = pres && (omp_target_is_present(&w[n / 2 - 1], omp_get_default_device()) != 0);
if (!shared_mem) pres = pres && (omp_target_is_present(&w[0], omp_get_default_device()) == 0);
if (!shared_mem) pres = pres && (omp_target_is_present(&w[n / 2], omp_get_default_device()) == 0);
#pragma omp target exit data map(iterator(k = 1:2), release: w[k:n / 2 - (k + 1) + 1])
bool t3_ = true;
for (int k4_ = 0; k4_ < n; k4_++)
t3_ = t3_ && (b[k4_] == ref[k4_]);
ok = t3_ && pres;
printf("RESULT ran_on_host=%d check=%s\n", host_flag, ok ? "pass" : "fail");
return 0;
}
// spec (requirement): An iterator modifier defines a set of iterators and is allowed on the map clause.
// "An iterator modifier is a unique, complex modifier that defines a set of iterators, each of which is an
// iterator-identifier and an associated set of values." (5.2 §3.2.6 p.67 l.6-8)
// spec (requirement): A list item in a map clause may use any iterator-identifier defined in the iterator modifier of
// that clause.
// "A list item in a map clause may reference any iterator-identifier defined in its iterator modifier." (5.2 §5.8.3
// p.151 l.13-14)
// spec (this program): target enter data and target exit data map w[k:n/2-(k+1)+1] for k = 1:2 with an iterator
// modifier; the compiler stops with an internal error instead of mapping the iterated sections or giving a
// diagnostic.
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
DeclRefExpr for Decl not entered in LocalDeclMap? UNREACHABLE executed at /workspace/clang/lib/CodeGen/CGExpr.cpp:3180! 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. cce-c-map-iterator-ice.c:7:5: Generating code for declaration 'main'
What the specification requires
The specification text this program relies on and the compiler is expected to support.
OpenMP 5.2 §3.2.6 p.67 l.6-8
An iterator modifier is a unique, complex modifier that defines a set of iterators, each of which is an iterator-identifier and an associated set of values.
OpenMP 5.2 §5.8.3 p.151 l.13-14
A list item in a map clause may reference any iterator-identifier defined in its iterator modifier.
Evidence
Same compiler, one change
- the same sections mapped without the iterator modifier (w[1:n/2 - 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 }; int main(void) { int a[n]; int b[n]; int ref[n]; int host_flag; bool ok; int k; int w[n]; bool pres; bool shared_mem; int zc[1]; for (int i = 1; i <= n; i++) { a[i - 1] = i; b[i - 1] = 0; ref[i - 1] = 2 * i + 1; } host_flag = 1; for (int k1_ = 0; k1_ < n; k1_++) w[k1_] = 0; #pragma omp target enter data map(to: w[1:n / 2 - 2 + 1]) pres = true; for (int k2_ = 0; k2_ < 1; k2_++) zc[k2_] = 0; #pragma omp target map(to: zc) { zc[0] = 1; } shared_mem = zc[0] == 1; #pragma omp target map(to: a) map(from: b) map(tofrom: host_flag) { host_flag = (omp_is_initial_device() ? 1 : 0); for (int i = 1; i <= n; i++) { b[i - 1] = 2 * a[i - 1] + 1; } } if (!shared_mem) pres = pres && (omp_target_is_present(&w[1], omp_get_default_device()) != 0); if (!shared_mem) pres = pres && (omp_target_is_present(&w[n / 2 - 1], omp_get_default_device()) != 0); if (!shared_mem) pres = pres && (omp_target_is_present(&w[0], omp_get_default_device()) == 0); if (!shared_mem) pres = pres && (omp_target_is_present(&w[n / 2], omp_get_default_device()) == 0); #pragma omp target exit data map(release: w[1:n / 2 - 2 + 1]) bool t3_ = true; for (int k4_ = 0; k4_ < n; k4_++) t3_ = t3_ && (b[k4_] == ref[k4_]); ok = t3_ && pres; 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, GCC 16.2.0
Reproducers of this issue
- clang-c-map-iterator-omp51
- fuzz/s1-6085-cce21
- catalog/cce-c-device-type-directive-name-ice
- catalog/cce-c-map-iterator-ice
- catalog/cce-c-motion-iterator-rejected
- catalog/cce-c-uses-allocators-traits-rejected
- catalog/cce21-c-declare-target-local-rejected
- catalog/cce-atomic-default-mem-order-acquire-release-rejected