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}} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
