Observed issues › afar-c-error-directive-in-target-link
amdclang -fopenmp-version=60: error at(compilation) severity(warning) in a target region fails at the device link (__kmpc_error undefined)
Legality: contested. An independent adjudicating reading finds the result not required: Both readings hold. The error directive with at(compilation) and severity(warning) displays an implementation-defined message; only fatal aborts compilation. The observed output is a warning, which is the specified action, and nothing restricts the directive inside a target region (it is pure when at is compilation). The observed result 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 | not supported | not supported | not supported | not supported | not supported | n/a | n/a |
| CCE 19.0.0 | correct | correct | correct | unknown not run | unknown not run | n/a | n/a |
| CCE 20.0.2 | correct | correct | correct | unknown not run | unknown not run | n/a | n/a |
| CCE 21.0.3 | correct | correct | correct | unknown not run | unknown not run | n/a | n/a |
| CCE 21.0.4 | correct | correct | correct | 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 |
|---|---|---|
| not supported | AFAR 24.3.0, MI300A unified AFAR 24.3.0, MI300A copy AFAR 24.3.0, MI250X AFAR 24.3.0, MI300X AFAR 24.3.0, MI350X | showafar-c-error-directive-in-target-link.c:9:48: warning: probe [-Wuser-defined-warnings]
9 | #pragma omp error at(compilation) severity(warning) message("probe")
| ^
1 warning generated.
afar-c-error-directive-in-target-link.c:9:48: warning: probe [-Wuser-defined-warnings]
9 | #pragma omp error at(compilation) severity(warning) message("probe")
| ^
1 warning generated. |
| fails another way | NVHPC 25.11, A100 separate NVHPC 25.11, A100 managed NVHPC 26.9, A100 separate NVHPC 26.9, A100 managed | show"afar-c-error-directive-in-target-link.c", line 9: error: invalid text in pragma
#pragma omp error at(compilation) severity(warning) message("probe")
^
1 error detected in the compilation of "afar-c-error-directive-in-target-link.c". |
Reproduce
AFAR 24.3.0, MI300A unified
$AFAR/bin/amdclang -fopenmp --offload-arch=gfx942 -fopenmp-version=60 -fopenmp-version=60 -O2 repro.c -o repro HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro
CCE 19.0.0, MI300A unified (runs correctly)
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=60 -O2 repro.c -o repro HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro
Program (32 lines)
// An error directive with at(compilation) severity(warning) in a
// target region (OpenMP 6.0 §10.1: at(compilation) acts at compile time
// only; a warning lets the compilation go on). Expected: the warning, a
// successful build, "RESULT ... check=pass".
// Observed (AFAR amdclang 24.0.0git, -fopenmp --offload-arch=gfx942
// -fopenmp-version=60 -O2, 2026-10-21): the warning is printed twice (host
// and device pass), then the device link fails: "ld.lld: error: undefined
// symbol: __kmpc_error >>> referenced by ...img.lto.o:(__omp_offloading_..
// _main_l..)": the device code calls the runtime error routine, which the
// device runtime lacks (amdclang lowers at(compilation) to a run-time
// call: placed in the host part of main the program builds, and its run
// prints "OMP: Warning #290: ... Encountered user-directed warning: probe",
// which at(compilation) should not). The Fortran form (amdflang) builds and
// runs; CCE craycc 19, 20, 21 build and run both placements. Makes the C message, severity and
// severity.warning programs (6.0) a link failure on AFAR.
#include <stdio.h>
#include <omp.h>
int main(void) {
int b[64];
int host_flag = 1;
#pragma omp target map(from: b) map(tofrom: host_flag)
{
#pragma omp error at(compilation) severity(warning) message("probe")
host_flag = omp_is_initial_device();
for (int i = 0; i < 64; i++) b[i] = 2 * i + 1;
}
int ok = 1;
for (int i = 0; i < 64; i++) ok = ok && b[i] == 2 * i + 1;
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
afar-c-error-directive-in-target-link.c:9:48: warning: probe [-Wuser-defined-warnings]
9 | #pragma omp error at(compilation) severity(warning) message("probe")
| ^
1 warning generated.
afar-c-error-directive-in-target-link.c:9:48: warning: probe [-Wuser-defined-warnings]
9 | #pragma omp error at(compilation) severity(warning) message("probe")
| ^
1 warning generated.What the specification requires
The specification text this program relies on and the compiler is expected to support.
OpenMP 6.0 §10.1
Evidence
Same compiler, one change
- the error directive removed: 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 b[64]; int host_flag = 1; #pragma omp target map(from: b) map(tofrom: host_flag) { host_flag = omp_is_initial_device(); for (int i = 0; i < 64; i++) b[i] = 2 * i + 1; } int ok = 1; for (int i = 0; i < 64; i++) ok = ok && b[i] == 2 * i + 1; printf("RESULT ran_on_host=%d check=%s\n", host_flag, ok ? "pass" : "fail"); return 0; }
Other compilers run it correctly
CCE 19.0.0, CCE 20.0.2, CCE 21.0.3, CCE 21.0.4, GCC 16.2.0