[Openmp-commits] [llvm] [openmp] [OpenMP][DeviceRTL] Implement __kmpc_error for the error directive (PR #220702)

Wed Sep 2 12:55:48 PDT 2026

llvmorg-github-actions[bot] wrote:




@llvm/pr-subscribers-offload

Author: Caroline Newcombe (cenewcombe)

<details>
<summary>Changes</summary>

### Description

`error` with `at(execution)` inside a `target` region compiled and generated a device image, but failed at device link with `undefined symbol: __kmpc_error`, since the entry point existed only in the host runtime. clang and flang both emit the call correctly, so this adds the missing device-side definition.

Behavior follows the host runtime and OpenMP 5.1, 2.5.4:

- `severity(warning)`: prints the message, execution continues.
- `severity(fatal)`: prints the message, execution aborts.

Tests added under `offload/test/offloading/`, covering both severities, an absent `message` clause, and C and Fortran.

### Deliberate omissions

- **Source location.** The host prints `file:line:column` from `ident_t`. The device ident carries the same information, but the DeviceRTL has no consumer of `psource` and the host parser (`Shared/SourceInfo.h`) uses `std::string`, so it is not usable there. Planned as a follow-up.
- **OMPT callback.** The spec requires an `ompt_callback_error` dispatch. The DeviceRTL has no OMPT support at all, so this is not implemented for device execution.
- **Message output requires GPU libc.** Without it, `vprintf` on AMDGPU is a stub and no message is printed, matching the existing behavior of device assertions and user `printf`. The abort/continue behavior is unaffected.
- **A fatal message can be lost when stdout is redirected**, because device `printf` goes to buffered stdout and the trap aborts before the buffer is flushed. This matches `__assert_fail_internal`.

### Notes
Part of #204240 
Assisted-by: Github Copilot

---
Full diff: https://github.com/llvm/llvm-project/pull/220702.diff


5 Files Affected:

- (added) offload/test/offloading/error_directive.c (+39) 
- (added) offload/test/offloading/fortran/error_directive.f90 (+24) 
- (added) offload/test/offloading/fortran/error_directive_fatal.f90 (+20) 
- (modified) openmp/device/include/Interface.h (+4) 
- (modified) openmp/device/src/Misc.cpp (+12) 


``````````diff

diff --git a/offload/test/offloading/error_directive.c b/offload/test/offloading/error_directive.c
new file mode 100644
index 0000000000000..bf8c8a7bf98be
--- /dev/null
+++ b/offload/test/offloading/error_directive.c
@@ -0,0 +1,39 @@
+// Test the `error` directive with `at(execution)` inside a target region.
+//
+// REQUIRES: libc
+
+// RUN: %libomptarget-compile-generic -fopenmp-version=51 && \
+// RUN:   %libomptarget-run-generic 2>&1 | %fcheck-generic
+// RUN: %libomptarget-compile-generic -fopenmp-version=51 -DFATAL && \
+// RUN:   %libomptarget-run-fail-generic
+
+#include <stdio.h>
+
+int main(void) {
+#ifdef FATAL
+#pragma omp target
+  {
+#pragma omp error at(execution) severity(fatal) message("fatal message")
+  }
+  printf("unreachable\n");
+#else
+#pragma omp target
+  {
+#pragma omp error at(execution) severity(warning) message("warning message")
+  }
+
+  // No MESSAGE clause, so the runtime receives a null message pointer.
+#pragma omp target
+  {
+#pragma omp error at(execution) severity(warning)
+  }
+#endif
+  return 0;
+}
+
+// Device output is flushed after host output, so host prints are not checked.
+// The fatal case checks only the exit status: its message is lost when the trap
+// aborts before the buffered stdout is flushed.
+
+// CHECK: user-directed warning: warning message.
+// CHECK: user-directed warning.
diff --git a/offload/test/offloading/fortran/error_directive.f90 b/offload/test/offloading/fortran/error_directive.f90
new file mode 100644
index 0000000000000..882d603e72536
--- /dev/null
+++ b/offload/test/offloading/fortran/error_directive.f90
@@ -0,0 +1,24 @@
+! Test the `error` directive with `at(execution)` inside a target region.
+!
+! REQUIRES: flang, libc
+
+! RUN: %libomptarget-compile-fortran-generic -fopenmp-version=51 && \
+! RUN:   %libomptarget-run-generic 2>&1 | %fcheck-generic
+
+program error_directive
+  implicit none
+
+  !$omp target
+  !$omp error at(execution) severity(warning) message("warning message")
+  !$omp end target
+
+  ! No MESSAGE clause, so the runtime receives a null message pointer.
+  !$omp target
+  !$omp error at(execution) severity(warning)
+  !$omp end target
+end program error_directive
+
+! Device output is flushed after host output, so host prints are not checked.
+
+! CHECK: user-directed warning: warning message.
+! CHECK: user-directed warning.
diff --git a/offload/test/offloading/fortran/error_directive_fatal.f90 b/offload/test/offloading/fortran/error_directive_fatal.f90
new file mode 100644
index 0000000000000..bef361192bcd1
--- /dev/null
+++ b/offload/test/offloading/fortran/error_directive_fatal.f90
@@ -0,0 +1,20 @@
+! Test `severity(fatal)` on the `error` directive with `at(execution)` inside a
+! target region: execution aborts.
+!
+! Only the exit status is checked: the message is lost when the trap aborts
+! before the buffered stdout is flushed. error_directive.f90 covers the text.
+!
+! REQUIRES: flang, libc
+
+! RUN: %libomptarget-compile-fortran-generic -fopenmp-version=51 && \
+! RUN:   %libomptarget-run-fail-generic
+
+program error_directive_fatal
+  implicit none
+
+  !$omp target
+  !$omp error at(execution) severity(fatal) message("fatal message")
+  !$omp end target
+
+  print *, "unreachable"
+end program error_directive_fatal
diff --git a/openmp/device/include/Interface.h b/openmp/device/include/Interface.h
index 601694871597b..98b9161408b46 100644
--- a/openmp/device/include/Interface.h
+++ b/openmp/device/include/Interface.h
@@ -357,6 +357,10 @@ void __kmpc_taskloop(IdentTy *Loc, uint32_t TId,
 int32_t __kmpc_cancellationpoint(IdentTy *Loc, int32_t TId, int32_t CancelVal);
 
 int32_t __kmpc_cancel(IdentTy *Loc, int32_t TId, int32_t CancelVal);
+
+/// Report a user-directed error. \p Severity matches kmp_severity_t: 1 is a
+/// warning and execution continues, 2 is fatal and execution aborts.
+void __kmpc_error(IdentTy *Loc, int32_t Severity, const char *Message);
 ///}
 
 /// Shuffle
diff --git a/openmp/device/src/Misc.cpp b/openmp/device/src/Misc.cpp
index f31639a46da18..3393d935bc356 100644
--- a/openmp/device/src/Misc.cpp
+++ b/openmp/device/src/Misc.cpp
@@ -77,6 +77,18 @@ int32_t __kmpc_cancellationpoint(IdentTy *, int32_t, int32_t) { return 0; }
 
 int32_t __kmpc_cancel(IdentTy *, int32_t, int32_t) { return 0; }
 
+// TODO: Report the source location from Loc->psource like the host runtime.
+void __kmpc_error(IdentTy *, int32_t Severity, const char *Message) {
+  const char *Kind = Severity == 1 ? "warning" : "error";
+  if (Message)
+    ompx::printf("OMP: Encountered user-directed %s: %s.\n", Kind, Message);
+  else
+    ompx::printf("OMP: Encountered user-directed %s.\n", Kind);
+
+  if (Severity != 1)
+    __builtin_trap();
+}
+
 double omp_get_wtick(void) {
   // The number of ticks per second for the AMDGPU clock varies by card and can
   // only be retrieved by querying the driver. We rely on the device environment

``````````

</details>


https://github.com/llvm/llvm-project/pull/220702
