https://github.com/akash-manna-sky updated https://github.com/llvm/llvm-project/pull/226691
>From 1a9aaafd8e29101e5db9fadac9527f09b4d56928 Mon Sep 17 00:00:00 2001 From: Akash Manna <[email protected]> Date: Sat, 26 Sep 2026 18:25:33 +0530 Subject: [PATCH 1/3] [clang][OpenMP] Fix crash on global reference in target region inside a namespace-scope lambda SemaOpenMP::isOpenMPCapturedDecl figures out whether a global must be captured for a target region by walking the function scope stack down to the innermost OpenMP captured region, and it stops as soon as it meets an ordinary function scope. When a directive is finished, the capture initializers of its outermost region are built after every region has been popped, while the directive is still on the DSA stack. Inside a function that walk simply ends at the function's scope. A lambda or block at namespace scope has no scope underneath it at all, so the walk ran off the end of the stack and tripped the assertion. The self-referential declaration in the report is incidental: plain `int &r = x; auto l = [] { #pragma omp target r = 1; };` crashes the same way, and so do lambdas in default arguments and in variable template instantiations. Running out of scopes now means the same thing as reaching a function scope: the variable is not captured from the current scope, so the walk returns null. That is exactly what the same lambda already gets when it sits inside a function, and it matches what the neighbouring walk for the CheckScopeInfo case already does when it finds no OpenMP region. Fixes #223397 --- clang/docs/ReleaseNotes.md | 1 + clang/lib/Sema/SemaOpenMP.cpp | 6 +- ...rget_global_ref_namespace_scope_lambda.cpp | 55 +++++++++++++++++++ 3 files changed, 61 insertions(+), 1 deletion(-) create mode 100644 clang/test/OpenMP/target_global_ref_namespace_scope_lambda.cpp diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md index f4a34a37aff52..4164ab1da17f2 100644 --- a/clang/docs/ReleaseNotes.md +++ b/clang/docs/ReleaseNotes.md @@ -544,6 +544,7 @@ features cannot lower the translation-unit ABI level; - Fixed a crash when declaring a member template within a local class inside an OpenMP region. (#GH216052) - 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 a crash when a reference variable with static storage duration, such as a global reference, was used in an OpenMP `target` region inside a lambda or block at namespace scope. (#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/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index 2e4d9f2f82f0b..6cbfe113902a3 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -2473,7 +2473,11 @@ VarDecl *SemaOpenMP::isOpenMPCapturedDecl(ValueDecl *D, bool CheckScopeInfo, break; } } - assert(CSI && "Failed to find CapturedRegionScopeInfo"); + // A lambda or block at namespace scope has no enclosing function scope, + // so the walk can run out of scopes once all captured regions of the + // directive have been left. + 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}} +}; >From 35cdf6c208c7c48cc5154d11bd2932539e17dfd0 Mon Sep 17 00:00:00 2001 From: Akash Manna <[email protected]> Date: Sun, 27 Sep 2026 22:16:39 +0530 Subject: [PATCH 2/3] [clang][OpenMP] Handle target regions without a parent function in CodeGen A target region inside a lambda or block at namespace scope has no parent function declaration, so emitCommonOMPTargetDirective asserted while computing the kernel's parent name. Use the name of the function being emitted instead. --- clang/docs/ReleaseNotes.md | 2 +- clang/lib/CodeGen/CGStmtOpenMP.cpp | 8 ++++--- clang/lib/Sema/SemaOpenMP.cpp | 4 +--- ...rget_global_ref_namespace_scope_lambda.cpp | 21 ++++++++++++++++++- 4 files changed, 27 insertions(+), 8 deletions(-) diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md index 291e4f364a986..9441b8672b6d2 100644 --- a/clang/docs/ReleaseNotes.md +++ b/clang/docs/ReleaseNotes.md @@ -554,7 +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 a crash when a reference variable with static storage duration, such as a global reference, was used in an OpenMP `target` region inside a lambda or block at namespace scope. (#GH223397) +- 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 70948022c0b44..479ba4c0e9a59 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -2488,9 +2488,7 @@ VarDecl *SemaOpenMP::isOpenMPCapturedDecl(ValueDecl *D, bool CheckScopeInfo, break; } } - // A lambda or block at namespace scope has no enclosing function scope, - // so the walk can run out of scopes once all captured regions of the - // directive have been left. + // Lambdas and blocks at namespace scope have no enclosing function scope. if (!CSI) return nullptr; SmallVector<OpenMPDirectiveKind, 4> Regions; diff --git a/clang/test/OpenMP/target_global_ref_namespace_scope_lambda.cpp b/clang/test/OpenMP/target_global_ref_namespace_scope_lambda.cpp index 32a7e6dfd74e4..68ffdf9975112 100644 --- a/clang/test/OpenMP/target_global_ref_namespace_scope_lambda.cpp +++ b/clang/test/OpenMP/target_global_ref_namespace_scope_lambda.cpp @@ -1,55 +1,74 @@ // RUN: %clang_cc1 -verify -fopenmp -fblocks -fsyntax-only %s // RUN: %clang_cc1 -verify -fopenmp-simd -fblocks -fsyntax-only %s +// RUN: %clang_cc1 -fopenmp -fblocks -DCODEGEN -triple x86_64-unknown-linux-gnu -emit-llvm %s -o - | FileCheck %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. +// block at namespace scope used to assert in SemaOpenMP::isOpenMPCapturedDecl, +// and such target regions have no parent function declaration in CodeGen. int x; int &ref = x; auto lambda = []() { +// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}_l[[#@LINE+1]]{{"?}}( #pragma omp target ref = 42; }; auto nested_lambda = []() { return []() { +// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}_l[[#@LINE+1]]{{"?}}( #pragma omp target ref = 42; }; }; auto combined_directive = []() { +// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}_l[[#@LINE+1]]{{"?}}( #pragma omp target teams ref = 42; }; auto static_local = []() { static int &local_ref = x; +// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}_l[[#@LINE+1]]{{"?}}( #pragma omp target local_ref = 42; }; void (^block)() = ^{ +// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}_l[[#@LINE+1]]{{"?}}( #pragma omp target ref = 42; }; void default_argument(int = []() { +// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}_l[[#@LINE+1]]{{"?}}( #pragma omp target ref = 42; return 0; }()); template <int N> int variable_template = []() { +// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}_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(); +} + +#ifndef CODEGEN // 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}} }; +#endif >From 15a369edb7412a6a6859f89823646aef5d27f239 Mon Sep 17 00:00:00 2001 From: Akash Manna <[email protected]> Date: Mon, 28 Sep 2026 10:10:02 +0530 Subject: [PATCH 3/3] [clang][OpenMP] Move codegen checks into a separate test Keep target_global_ref_namespace_scope_lambda.cpp a Sema-only test and check the outlined target region names in a new target_global_ref_namespace_scope_lambda_codegen.cpp. --- ...rget_global_ref_namespace_scope_lambda.cpp | 21 +----- ...bal_ref_namespace_scope_lambda_codegen.cpp | 67 +++++++++++++++++++ 2 files changed, 68 insertions(+), 20 deletions(-) create mode 100644 clang/test/OpenMP/target_global_ref_namespace_scope_lambda_codegen.cpp diff --git a/clang/test/OpenMP/target_global_ref_namespace_scope_lambda.cpp b/clang/test/OpenMP/target_global_ref_namespace_scope_lambda.cpp index 68ffdf9975112..32a7e6dfd74e4 100644 --- a/clang/test/OpenMP/target_global_ref_namespace_scope_lambda.cpp +++ b/clang/test/OpenMP/target_global_ref_namespace_scope_lambda.cpp @@ -1,74 +1,55 @@ // RUN: %clang_cc1 -verify -fopenmp -fblocks -fsyntax-only %s // RUN: %clang_cc1 -verify -fopenmp-simd -fblocks -fsyntax-only %s -// RUN: %clang_cc1 -fopenmp -fblocks -DCODEGEN -triple x86_64-unknown-linux-gnu -emit-llvm %s -o - | FileCheck %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, -// and such target regions have no parent function declaration in CodeGen. +// block at namespace scope used to assert in SemaOpenMP::isOpenMPCapturedDecl. int x; int &ref = x; auto lambda = []() { -// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}_l[[#@LINE+1]]{{"?}}( #pragma omp target ref = 42; }; auto nested_lambda = []() { return []() { -// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}_l[[#@LINE+1]]{{"?}}( #pragma omp target ref = 42; }; }; auto combined_directive = []() { -// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}_l[[#@LINE+1]]{{"?}}( #pragma omp target teams ref = 42; }; auto static_local = []() { static int &local_ref = x; -// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}_l[[#@LINE+1]]{{"?}}( #pragma omp target local_ref = 42; }; void (^block)() = ^{ -// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}_l[[#@LINE+1]]{{"?}}( #pragma omp target ref = 42; }; void default_argument(int = []() { -// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}_l[[#@LINE+1]]{{"?}}( #pragma omp target ref = 42; return 0; }()); template <int N> int variable_template = []() { -// CHECK-DAG: define internal void @{{"?}}__omp_offloading_{{.*}}_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(); -} - -#ifndef CODEGEN // 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}} }; -#endif 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
