[clang] 2bf008c - [clang][OpenMP] Fix crashes on target regions inside namespace-scope lambdas and blocks (#226691)
via cfe-commits
cfe-commits at lists.llvm.org
Tue Sep 29 06:51:21 PDT 2026
Author: Akash Manna
Date: 2026-09-29T19:21:15+05:30
New Revision: 2bf008c3ad7fc34de4a6c9eb1bd8910457d48a34
URL: https://github.com/llvm/llvm-project/commit/2bf008c3ad7fc34de4a6c9eb1bd8910457d48a34
DIFF: https://github.com/llvm/llvm-project/commit/2bf008c3ad7fc34de4a6c9eb1bd8910457d48a34.diff
LOG: [clang][OpenMP] Fix crashes on target regions inside namespace-scope lambdas and blocks (#226691)
Fixes #223397
A `target` region inside a lambda or block at namespace scope crashed
clang in two places. In Sema, `isOpenMPCapturedDecl` decides whether a
global must be captured by walking the function scope stack down to the
innermost OpenMP captured region, stopping at an ordinary function
scope. The capture initializers of a directive's outermost region are
built after all of its regions have been popped. Inside a function that
walk ends at the function's scope, but a namespace-scope lambda or block
has nothing underneath it, so the walk ran off the stack and asserted.
This happens for any global reference, such as `int &r = x; auto l = []
{ #pragma omp target r = 1; };`. The self-referential declaration in the
report is incidental. Once past Sema, CodeGen asserted too: it names the
outlined kernel after the region's parent function, and such a region
has none, even without a reference.
In Sema, running out of scopes now means the same as reaching a function
scope, so the variable is not captured from the current scope. That is
what the same lambda already gets inside a function. In CodeGen, a
region with no parent function is named after the function being
emitted.
Added:
clang/test/OpenMP/target_global_ref_namespace_scope_lambda.cpp
clang/test/OpenMP/target_global_ref_namespace_scope_lambda_codegen.cpp
Modified:
clang/docs/ReleaseNotes.md
clang/lib/CodeGen/CGStmtOpenMP.cpp
clang/lib/Sema/SemaOpenMP.cpp
Removed:
################################################################################
diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md
index 49d669db5749c..7c204eb2ca577 100644
--- a/clang/docs/ReleaseNotes.md
+++ b/clang/docs/ReleaseNotes.md
@@ -554,6 +554,7 @@ features cannot lower the translation-unit ABI level;
(such as `std::map`) whose mapper lookup instantiates a class template specialization. (#GH154704)
- Fixed a bug where repeated #imports of modular headers in non-modular compilation were translated to #pragma clang module import. (#GH216924)
- Fixed an assertion when `#pragma omp declare simd` or `#pragma omp declare variant` is followed by another OpenMP declarative directive containing a qualified identifier. (#GH217204)
+- Fixed crashes on an OpenMP `target` region inside a lambda or block at namespace scope, including when the region used a global reference. (#GH223397)
- Fixed a crash when an `asm` label names the register for a global variable of incomplete type. (#GH219746)
- Fixed an ICE hat occurred when using `__imag int/float` as lvalue in assignment. (#GH119498)
- Fixed an assertion failure in `-Wsign-compare` when a negated or complemented vector of unsigned integers was compared against a signed constant. (#GH203575)
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index e4751a90d30b0..b5e4b01109cb3 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -7542,11 +7542,13 @@ static void emitCommonOMPTargetDirective(CodeGenFunction &CGF,
CGM.getDiags().Report(diag::err_missing_mandatory_offloading);
}
- assert(CGF.CurFuncDecl && "No parent declaration for target region!");
StringRef ParentName;
// In case we have Ctors/Dtors we use the complete type variant to produce
- // the mangling of the device outlined kernel.
- if (const auto *D = dyn_cast<CXXConstructorDecl>(CGF.CurFuncDecl))
+ // the mangling of the device outlined kernel. Lambdas and blocks at
+ // namespace scope have no parent function.
+ if (!CGF.CurFuncDecl)
+ ParentName = CGF.CurFn->getName();
+ else if (const auto *D = dyn_cast<CXXConstructorDecl>(CGF.CurFuncDecl))
ParentName = CGM.getMangledName(GlobalDecl(D, Ctor_Complete));
else if (const auto *D = dyn_cast<CXXDestructorDecl>(CGF.CurFuncDecl))
ParentName = CGM.getMangledName(GlobalDecl(D, Dtor_Complete));
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 41cfc1dcd3649..479ba4c0e9a59 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -2488,7 +2488,9 @@ VarDecl *SemaOpenMP::isOpenMPCapturedDecl(ValueDecl *D, bool CheckScopeInfo,
break;
}
}
- assert(CSI && "Failed to find CapturedRegionScopeInfo");
+ // Lambdas and blocks at namespace scope have no enclosing function scope.
+ if (!CSI)
+ return nullptr;
SmallVector<OpenMPDirectiveKind, 4> Regions;
getOpenMPCaptureRegions(Regions,
DSAStack->getDirective(CSI->OpenMPLevel));
diff --git a/clang/test/OpenMP/target_global_ref_namespace_scope_lambda.cpp b/clang/test/OpenMP/target_global_ref_namespace_scope_lambda.cpp
new file mode 100644
index 0000000000000..32a7e6dfd74e4
--- /dev/null
+++ b/clang/test/OpenMP/target_global_ref_namespace_scope_lambda.cpp
@@ -0,0 +1,55 @@
+// RUN: %clang_cc1 -verify -fopenmp -fblocks -fsyntax-only %s
+// RUN: %clang_cc1 -verify -fopenmp-simd -fblocks -fsyntax-only %s
+
+// A reference without local storage used in a target region inside a lambda or
+// block at namespace scope used to assert in SemaOpenMP::isOpenMPCapturedDecl.
+
+int x;
+int &ref = x;
+
+auto lambda = []() {
+#pragma omp target
+ ref = 42;
+};
+
+auto nested_lambda = []() {
+ return []() {
+#pragma omp target
+ ref = 42;
+ };
+};
+
+auto combined_directive = []() {
+#pragma omp target teams
+ ref = 42;
+};
+
+auto static_local = []() {
+ static int &local_ref = x;
+#pragma omp target
+ local_ref = 42;
+};
+
+void (^block)() = ^{
+#pragma omp target
+ ref = 42;
+};
+
+void default_argument(int = []() {
+#pragma omp target
+ ref = 42;
+ return 0;
+}());
+
+template <int N> int variable_template = []() {
+#pragma omp target
+ ref = N;
+ return 0;
+}();
+int instantiation = variable_template<1>;
+
+// Reproducer from GH223397.
+int &foo = []() { // expected-error {{non-const lvalue reference to type 'int' cannot bind to a temporary of type '(lambda at}}
+#pragma omp target
+ foo(42); // expected-error {{called object type 'int' is not a function or function pointer}}
+};
diff --git a/clang/test/OpenMP/target_global_ref_namespace_scope_lambda_codegen.cpp b/clang/test/OpenMP/target_global_ref_namespace_scope_lambda_codegen.cpp
new file mode 100644
index 0000000000000..3ac1f55ca7a96
--- /dev/null
+++ b/clang/test/OpenMP/target_global_ref_namespace_scope_lambda_codegen.cpp
@@ -0,0 +1,67 @@
+// RUN: %clang_cc1 -verify -fopenmp -fblocks -triple x86_64-unknown-linux-gnu -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp-simd -fblocks -triple x86_64-unknown-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=SIMD-ONLY
+// expected-no-diagnostics
+
+// Target regions inside lambdas and blocks at namespace scope have no parent
+// function declaration, so they are named after the function being emitted.
+
+int x;
+int &ref = x;
+
+auto lambda = []() {
+// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}clEv_l[[#@LINE+1]]{{"?}}(
+#pragma omp target
+ ref = 42;
+};
+
+auto nested_lambda = []() {
+ return []() {
+// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}clEv_l[[#@LINE+1]]{{"?}}(
+#pragma omp target
+ ref = 42;
+ };
+};
+
+auto combined_directive = []() {
+// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}clEv_l[[#@LINE+1]]{{"?}}(
+#pragma omp target teams
+ ref = 42;
+};
+
+auto static_local = []() {
+ static int &local_ref = x;
+// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}clEv_l[[#@LINE+1]]{{"?}}(
+#pragma omp target
+ local_ref = 42;
+};
+
+void (^block)() = ^{
+// CHECK-DAG: define internal void @__omp_offloading_{{.*}}_block_block_invoke_l[[#@LINE+1]](
+#pragma omp target
+ ref = 42;
+};
+
+void default_argument(int = []() {
+// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}clEv_l[[#@LINE+1]]{{"?}}(
+#pragma omp target
+ ref = 42;
+ return 0;
+}());
+
+template <int N> int variable_template = []() {
+// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}clEv_l[[#@LINE+1]]{{"?}}(
+#pragma omp target
+ ref = N;
+ return 0;
+}();
+int instantiation = variable_template<1>;
+
+void use() {
+ lambda();
+ nested_lambda()();
+ combined_directive();
+ static_local();
+ default_argument();
+}
+
+// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
More information about the cfe-commits
mailing list