Observed issues › cce-c-openmp45-device-link
Appears not to supportRefused at link time · C · OpenMP · Open
craycc -fopenmp-version=45: omp_is_initial_device called in a target region is an undefined symbol at the device link
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 | correct | correct | correct | n/a | n/a |
| CCE 19.0.0 | not supported | not supported | not supported | unknown not run | unknown not run | n/a | n/a |
| CCE 20.0.2 | not supported | not supported | not supported | unknown not run | unknown not run | n/a | n/a |
| CCE 21.0.3 | not supported | not supported | not supported | unknown not run | unknown not run | n/a | n/a |
| CCE 21.0.4 | not supported | not supported | not supported | 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 |
|---|---|---|
| not supported | CCE 19.0.0, MI300A unified CCE 19.0.0, MI300A copy | showlld: error: undefined symbol: omp_is_initial_device >>> referenced by <dir>/cooltmpdir-26TMiI/repro-cce-openmp__llc.amdgpu:(omp_is_initial_device_cce$omp_redun_clone) >>> referenced by <dir>/cooltmpdir-26TMiI/repro-cce-openmp__llc.amdgpu:(omp_is_initial_device_cce$omp_redun_clone) clang: error: linker command failed with exit code 1 (use -v to see invocation) |
| not supported | CCE 19.0.0, MI250X | showlld: error: undefined symbol: omp_is_initial_device >>> referenced by <dir>/cooltmpdir-BV8iAi/repro-cce-openmp__llc.amdgpu:(omp_is_initial_device_cce$omp_redun_clone) >>> referenced by <dir>/cooltmpdir-BV8iAi/repro-cce-openmp__llc.amdgpu:(omp_is_initial_device_cce$omp_redun_clone) clang: error: linker command failed with exit code 1 (use -v to see invocation) |
| not supported | CCE 20.0.2, MI300A unified CCE 20.0.2, MI300A copy | showlld: error: undefined symbol: omp_is_initial_device >>> referenced by <dir>/cooltmpdir-L7jmND/repro-cce-openmppost_lld.amdgpu.lto.o:(__omp_offloading_14_1c76d3f0_main_l5) >>> referenced by <dir>/cooltmpdir-L7jmND/repro-cce-openmppost_lld.amdgpu.lto.o:(__omp_offloading_14_1c76d3f0_main_l5) clang: error: linker command failed with exit code 1 (use -v to see invocation) |
| not supported | CCE 20.0.2, MI250X | showlld: error: undefined symbol: omp_is_initial_device >>> referenced by <dir>/cooltmpdir-M6NDdt/repro-cce-openmppost_lld.amdgpu.lto.o:(__omp_offloading_14_378d8051_main_l5) >>> referenced by <dir>/cooltmpdir-M6NDdt/repro-cce-openmppost_lld.amdgpu.lto.o:(__omp_offloading_14_378d8051_main_l5) clang: error: linker command failed with exit code 1 (use -v to see invocation) |
| not supported | CCE 21.0.3, MI300A unified CCE 21.0.3, MI300A copy | showlld: error: undefined symbol: omp_is_initial_device >>> referenced by <dir>/cooltmpdir-982PZp/repro-cce-openmppost_lld.amdgpu.lto.o:(__omp_offloading_14_1c7336fd_main_l5) >>> referenced by <dir>/cooltmpdir-982PZp/repro-cce-openmppost_lld.amdgpu.lto.o:(__omp_offloading_14_1c7336fd_main_l5) craycc: error: linker command failed with exit code 1 (use -v to see invocation) |
| not supported | CCE 21.0.3, MI250X | showlld: error: undefined symbol: omp_is_initial_device >>> referenced by <dir>/cooltmpdir-EDa0sF/repro-cce-openmppost_lld.amdgpu.lto.o:(__omp_offloading_14_3792f710_main_l5) >>> referenced by <dir>/cooltmpdir-EDa0sF/repro-cce-openmppost_lld.amdgpu.lto.o:(__omp_offloading_14_3792f710_main_l5) craycc: error: linker command failed with exit code 1 (use -v to see invocation) |
| not supported | CCE 21.0.4, MI300A unified CCE 21.0.4, MI300A copy | showlld: error: undefined symbol: omp_is_initial_device >>> referenced by <dir>/cooltmpdir-Cezbi8/repro-cce-openmppost_lld.amdgpu.lto.o:(__omp_offloading_14_1c761367_main_l5) >>> referenced by <dir>/cooltmpdir-Cezbi8/repro-cce-openmppost_lld.amdgpu.lto.o:(__omp_offloading_14_1c761367_main_l5) craycc: error: linker command failed with exit code 1 (use -v to see invocation) |
| not supported | CCE 21.0.4, MI250X | showlld: error: undefined symbol: omp_is_initial_device >>> referenced by <dir>/cooltmpdir-16FYZP/repro-cce-openmppost_lld.amdgpu.lto.o:(__omp_offloading_14_3790e932_main_l5) >>> referenced by <dir>/cooltmpdir-16FYZP/repro-cce-openmppost_lld.amdgpu.lto.o:(__omp_offloading_14_3790e932_main_l5) craycc: error: linker command failed with exit code 1 (use -v to see invocation) |
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=45 -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=45 -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=45 -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=45 -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=45 -fopenmp-version=45 -O2 repro.c -o repro HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro
Program (29 lines)
/* Finding (C): with -fopenmp-version=45 CCE cannot link
omp_is_initial_device called in a target region.
Legal OpenMP 4.5 C (4.5 §3.2.34: omp_is_initial_device; runtime
library routines may be called in target regions).
Expected: "on_host=0".
Observed: CCE craycc 19.0.0, 20.0.2 and 21.0.3 with
-fopenmp-version=45 fail to link:
lld: error: undefined symbol: omp_is_initial_device
>>> referenced by ...-cce-openmp__llc.amdgpu:
(omp_is_initial_device_cce$omp_redun_clone)
With no version flag (default 5.1; CCE 19 also at 50 and 52) it
links and prints "on_host=0". The first omp_vv C baseline (each test
at the version of its spec directory) lost 115 of 315 tests on CCE to
this; the baseline now compiles 4.5-5.1 tests with the default.
craycc -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa
-Xopenmp-target=amdgcn-amd-amdhsa -march=gfx942
-fopenmp-version=45 -O2 */
#include <stdio.h>
#include <omp.h>
int main(void) {
int on_host = 1;
#pragma omp target map(tofrom: on_host)
on_host = omp_is_initial_device();
printf("on_host=%d\n", on_host);
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
lld: error: undefined symbol: omp_is_initial_device >>> referenced by <dir>/cooltmpdir-26TMiI/repro-cce-openmp__llc.amdgpu:(omp_is_initial_device_cce$omp_redun_clone) >>> referenced by <dir>/cooltmpdir-26TMiI/repro-cce-openmp__llc.amdgpu:(omp_is_initial_device_cce$omp_redun_clone) clang: error: linker command failed with exit code 1 (use -v to see invocation)
What the specification requires
The specification text this program relies on and the compiler is expected to support.
OpenMP 4.5 §3.2.34
Evidence
Same compiler, one change
- built without -fopenmp-version=45 (the default, 5.1): 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) { int on_host = 1; #pragma omp target map(tofrom: on_host) on_host = omp_is_initial_device(); printf("on_host=%d\n", on_host); return 0; }
Other compilers run it correctly
AFAR 24.3.0, NVHPC 25.11, NVHPC 26.9, GCC 16.2.0