[flang-commits] [flang] [llvm] [openmp] [OpenMP][DeviceRTL] Implement __kmpc_error for the error directive (PR #220702)
Caroline Newcombe via flang-commits
flang-commits at lists.llvm.org
Mon Sep 14 13:04:26 PDT 2026
https://github.com/cenewcombe updated https://github.com/llvm/llvm-project/pull/220702
>From 07ec4ceb6d743a8b242b958de59bdf8c964555a4 Mon Sep 17 00:00:00 2001
From: Caroline Newcombe <caroline.newcombe at hpe.com>
Date: Wed, 2 Sep 2026 11:45:35 -0500
Subject: [PATCH 1/3] [OpenMP][DeviceRTL] Implement __kmpc_error for the error
directive
---
offload/test/offloading/error_directive.c | 39 +++++++++++++++++++
.../offloading/fortran/error_directive.f90 | 24 ++++++++++++
.../fortran/error_directive_fatal.f90 | 20 ++++++++++
openmp/device/include/Interface.h | 4 ++
openmp/device/src/Misc.cpp | 12 ++++++
5 files changed, 99 insertions(+)
create mode 100644 offload/test/offloading/error_directive.c
create mode 100644 offload/test/offloading/fortran/error_directive.f90
create mode 100644 offload/test/offloading/fortran/error_directive_fatal.f90
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
>From 99502758dbe027599638ba45aea766061622e8e9 Mon Sep 17 00:00:00 2001
From: Caroline Newcombe <caroline.newcombe at hpe.com>
Date: Mon, 14 Sep 2026 15:03:25 -0500
Subject: [PATCH 2/3] Mark Fortran error directive tests unsupported for LTO
---
offload/test/offloading/fortran/error_directive.f90 | 3 +++
offload/test/offloading/fortran/error_directive_fatal.f90 | 3 +++
2 files changed, 6 insertions(+)
diff --git a/offload/test/offloading/fortran/error_directive.f90 b/offload/test/offloading/fortran/error_directive.f90
index 882d603e72536..2b2ece3d42ffc 100644
--- a/offload/test/offloading/fortran/error_directive.f90
+++ b/offload/test/offloading/fortran/error_directive.f90
@@ -2,6 +2,9 @@
!
! REQUIRES: flang, libc
+! flang does not accept -foffload-lto.
+! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+
! RUN: %libomptarget-compile-fortran-generic -fopenmp-version=51 && \
! RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic
diff --git a/offload/test/offloading/fortran/error_directive_fatal.f90 b/offload/test/offloading/fortran/error_directive_fatal.f90
index bef361192bcd1..b27e0e2a67e1a 100644
--- a/offload/test/offloading/fortran/error_directive_fatal.f90
+++ b/offload/test/offloading/fortran/error_directive_fatal.f90
@@ -6,6 +6,9 @@
!
! REQUIRES: flang, libc
+! flang does not accept -foffload-lto.
+! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+
! RUN: %libomptarget-compile-fortran-generic -fopenmp-version=51 && \
! RUN: %libomptarget-run-fail-generic
>From 8a643fe0e0332d09b907960a37eda7e15f1e9208 Mon Sep 17 00:00:00 2001
From: Caroline Newcombe <caroline.newcombe at hpe.com>
Date: Mon, 14 Sep 2026 15:04:09 -0500
Subject: [PATCH 3/3] Update OpenMP support documentation for the error
directive
---
flang/docs/OpenMPSupport.md | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/flang/docs/OpenMPSupport.md b/flang/docs/OpenMPSupport.md
index 5a5285421e724..75cbf7291eeba 100644
--- a/flang/docs/OpenMPSupport.md
+++ b/flang/docs/OpenMPSupport.md
@@ -198,7 +198,7 @@ Parser/Semantics, MLIR, Lowering, or the OpenMPIRBuilder.
| target_device selector | <span class="part">partial</span> | | Semantics coverage exists for target_device selectors in metadirective/declare-variant matching (`flang/test/Semantics/OpenMP/metadirective-device.f90`, `flang/test/Semantics/OpenMP/declare-variant-match.f90`). | [llvm/llvm-project#123243](https://github.com/llvm/llvm-project/pull/123243), [llvm/llvm-project#206714](https://github.com/llvm/llvm-project/pull/206714) |
| adjust_args and append_args on declare variant | <span class="none">unclaimed</span> | | Parsing accepts forms, but semantics currently diagnose both clauses as not yet implemented (`flang/test/Semantics/OpenMP/declare-variant-match.f90`, `flang/test/Semantics/OpenMP/declare-variant.f90`). | [llvm/llvm-project#206714](https://github.com/llvm/llvm-project/pull/206714) |
| indirect clause on declare target | <span class="part">partial</span> | | Parser coverage exists (`flang/test/Parser/OpenMP/declare-target-indirect-tree.f90`), while lowering remains TODO-tracked (`flang/test/Lower/OpenMP/Todo/omp-clause-indirect.f90`). | [llvm/llvm-project#143505](https://github.com/llvm/llvm-project/pull/143505) |
-| error directive | <span class="part">partial</span> | cenewcombe | Semantics coverage exists (`flang/test/Semantics/OpenMP/error.f90`), but lowering still has explicit TODO coverage (`flang/test/Lower/OpenMP/Todo/error.f90`). | [llvm/llvm-project#121509](https://github.com/llvm/llvm-project/pull/121509), [llvm/llvm-project#206175](https://github.com/llvm/llvm-project/pull/206175) |
+| error directive | <span class="good">done</span> | cenewcombe | Semantics, lowering, and runtime support are complete. `at(compilation)` is diagnosed at compile time and `at(execution)` is reported at run time, on both the host and inside `target` regions. | [llvm/llvm-project#121509](https://github.com/llvm/llvm-project/pull/121509), [llvm/llvm-project#206175](https://github.com/llvm/llvm-project/pull/206175), [llvm/llvm-project#220702](https://github.com/llvm/llvm-project/pull/220702) |
| nothing directive | <span class="good">done</span> | | Parser and lowering coverage exists for standalone and metadirective-selected forms (`flang/test/Parser/OpenMP/nothing.f90`, `flang/test/Lower/OpenMP/nothing.f90`, `flang/test/Lower/OpenMP/metadirective-nothing.f90`). | [llvm/llvm-project#193664](https://github.com/llvm/llvm-project/pull/193664), [llvm/llvm-project#202679](https://github.com/llvm/llvm-project/pull/202679) |
| tile and unroll constructs | <span class="part">partial</span> | | Semantics coverage exists across tile/unroll and loop-transformation tests (`flang/test/Semantics/OpenMP/tile01.f90`, `flang/test/Semantics/OpenMP/tile09.f90`, `flang/test/Semantics/OpenMP/loop-transformation-construct01.f90`), with additional lowering/transform completeness work ongoing. The `full` and `partial` clauses on `unroll` are lowered (`flang/test/Lower/OpenMP/unroll-full01.f90`, `flang/test/Lower/OpenMP/unroll-partial01.f90`). | [llvm/llvm-project#160298](https://github.com/llvm/llvm-project/pull/160298), [llvm/llvm-project#185296](https://github.com/llvm/llvm-project/pull/185296), [llvm/llvm-project#188025](https://github.com/llvm/llvm-project/pull/188025) |
| scope construct | <span class="part">partial</span> | | Scope construct support is available, with follow-on completeness work still in progress for some combinations (see also OpenMP 5.2 scope-related rows). | [llvm/llvm-project#113700](https://github.com/llvm/llvm-project/pull/113700), [llvm/llvm-project#193098](https://github.com/llvm/llvm-project/pull/193098) |
More information about the flang-commits
mailing list