Author: Akash Manna Date: 2026-09-25T06:46:28-04:00 New Revision: e34e31190233a5733d6b7751160a29e50cd8b41c
URL: https://github.com/llvm/llvm-project/commit/e34e31190233a5733d6b7751160a29e50cd8b41c DIFF: https://github.com/llvm/llvm-project/commit/e34e31190233a5733d6b7751160a29e50cd8b41c.diff LOG: [clang][OpenMP] Cast by-reference captures to the address space of the captured field (#226229) Fixes #140069 Sema strips all qualifiers, including the address space, from the type of a variable captured by an OpenMP region, so the outlined function for a `target` region takes `int __seg_gs a;` as a plain `int &`, i.e. a pointer in the default address space. That is the intended model: inside the region the variable is an ordinary object, and the host's segment address space means nothing on a device. `GenerateOpenMPCapturedVars` ignored the captured field's type for by-reference captures though, and passed the variable's own address. With `map(alloc: a)` that is a `ptr addrspace(256)` argument for a `ptr` parameter, and emitting the host fallback call tripped the "Calling a function with a bad signature" assertion. The same call is emitted on the `omp_offload.failed` path when offloading targets are given, so this wasn't limited to host-only compiles. By-reference captures now convert the address to the type of the captured field, adding an `addrspacecast` only when the address spaces differ. On x86 the cast doesn't change the pointer value, so the region gets the linear address of the object, which is also what the mapping passes to the runtime and what GCC does for the same code. Captures whose address already has the right type are untouched. Added: clang/test/OpenMP/target_map_address_space_codegen.c Modified: clang/docs/ReleaseNotes.md clang/lib/CodeGen/CGStmtOpenMP.cpp Removed: ################################################################################ diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md index f6cdca91cea10..702ff4a17d5c0 100644 --- a/clang/docs/ReleaseNotes.md +++ b/clang/docs/ReleaseNotes.md @@ -540,6 +540,7 @@ features cannot lower the translation-unit ABI level; - Fixed an assertion caused by Microsoft integer literals exceeding the maximum value. (#GH212504) - Fixed an assertion failure when a value of a Unicode character type (`char8_t`, `char16_t`, `char32_t`) was implicitly splatted to a vector of the same element type, e.g. when comparing an `ext_vector_type` of `char32_t` with one of its elements. (#GH202317) - Fixed a crash when checking scalar type with excess braces. (#GH69213), (#GH137845), (#GH198767), (#GH207566), (#GH106180) +- Fixed an assertion failure when a global variable in a non-default address space, such as one declared with `__seg_gs`, is mapped into an OpenMP `target` region. (#GH140069) - Fixed an assertion crash when instantiating a nested requirement with an invalid constraint. (#GH213575) - Clang now defines the GCC-compatible predefined macro `__SIG_ATOMIC_TYPE__`. (#GH213895) - Fixed IEEE f128 complex mul/div using the IBM f128 libcalls on powerpc. (#GH216820) diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp index 33ded363f948e..e4751a90d30b0 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -493,7 +493,12 @@ void CodeGenFunction::GenerateOpenMPCapturedVars( CapturedVars.push_back(CV); } else { assert(CurCap->capturesVariable() && "Expected capture by reference."); - CapturedVars.push_back(EmitLValue(*I).getAddress().emitRawPointer(*this)); + llvm::Value *Addr = EmitLValue(*I).getAddress().emitRawPointer(*this); + // Sema strips the address space from the type of the captured field. + llvm::Type *ArgTy = ConvertType(CurField->getType()); + if (Addr->getType() != ArgTy) + Addr = performAddrSpaceCast(Addr, ArgTy); + CapturedVars.push_back(Addr); } } } diff --git a/clang/test/OpenMP/target_map_address_space_codegen.c b/clang/test/OpenMP/target_map_address_space_codegen.c new file mode 100644 index 0000000000000..36ee8f59aeb28 --- /dev/null +++ b/clang/test/OpenMP/target_map_address_space_codegen.c @@ -0,0 +1,41 @@ +// RUN: %clang_cc1 -verify -fopenmp -triple x86_64-unknown-linux-gnu -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -verify -fopenmp -triple x86_64-unknown-linux-gnu -fopenmp-targets=x86_64-unknown-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,OFFLOAD +// expected-no-diagnostics + +// The outlined target region takes a variable from a non-default address space +// as a plain pointer, so the host fallback call has to cast its address (GH140069). + +int __seg_gs a; +int b; + +// CHECK-DAG: @a = {{.*}}addrspace(256) global i32 0 +// CHECK-DAG: @b = {{.*}}global i32 0 + +// CHECK-LABEL: define {{.*}}void @f( +// OFFLOAD: call i32 @__tgt_target_kernel( +// CHECK: call void @[[OUTLINED:__omp_offloading_[0-9a-z]+_[0-9a-z]+_f_l[0-9]+]](ptr addrspacecast (ptr addrspace(256) @a to ptr), ptr null) +void f(void) { +#pragma omp target map(alloc: a) map(from: b) + { + a = 0; + } +} + +// CHECK: define internal void @[[OUTLINED]](ptr noundef nonnull align 4 dereferenceable(4) %{{.+}}, ptr noalias noundef %{{.+}}) +// CHECK-NOT: addrspace(256) +// CHECK: store i32 0, ptr %{{.+}}, align 4 + +// CHECK-LABEL: define {{.*}}void @g( +// OFFLOAD: call i32 @__tgt_target_kernel( +// CHECK: call void @[[OUTLINED2:__omp_offloading_[0-9a-z]+_[0-9a-z]+_g_l[0-9]+]](ptr addrspacecast (ptr addrspace(256) @a to ptr), ptr @b, ptr null) +void g(void) { +#pragma omp target map(alloc: a) map(from: b) + { + a = 321; + b = a; + } +} + +// CHECK: define internal void @[[OUTLINED2]](ptr noundef nonnull align 4 dereferenceable(4) %{{.+}}, ptr noundef nonnull align 4 dereferenceable(4) %{{.+}}, ptr noalias noundef %{{.+}}) +// CHECK-NOT: addrspace(256) +// CHECK: store i32 321, ptr %{{.+}}, align 4 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
