https://github.com/akash-manna-sky created https://github.com/llvm/llvm-project/pull/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. >From fd2180c8347c3bfc8e2bab4187c8bc98bb1bd5ef Mon Sep 17 00:00:00 2001 From: Akash Manna <[email protected]> Date: Thu, 24 Sep 2026 22:19:12 +0530 Subject: [PATCH] [clang][OpenMP] Cast by-reference captures to the address space of the captured field Sema strips the qualifiers, including the address space, from the type of a variable captured by an OpenMP region, so the outlined function takes a __seg_gs global mapped into a target region as a pointer in the default address space. GenerateOpenMPCapturedVars still passed the variable's own address, so the host fallback call was built with a ptr addrspace(256) argument for a ptr parameter and hit the "Calling a function with a bad signature" assertion. Convert the address of a by-reference capture to the type of the captured field when the two differ. Fixes #140069 --- clang/docs/ReleaseNotes.md | 1 + clang/lib/CodeGen/CGStmtOpenMP.cpp | 7 +++- .../OpenMP/target_map_address_space_codegen.c | 41 +++++++++++++++++++ 3 files changed, 48 insertions(+), 1 deletion(-) create mode 100644 clang/test/OpenMP/target_map_address_space_codegen.c diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md index f4a34a37aff52e..5de37bb503c840 100644 --- a/clang/docs/ReleaseNotes.md +++ b/clang/docs/ReleaseNotes.md @@ -547,6 +547,7 @@ features cannot lower the translation-unit ABI level; - 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) +- 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) #### Bug Fixes to Compiler Builtins diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp index 7c3b30c6cedc0e..aca738c03a12e5 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -463,7 +463,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 00000000000000..36ee8f59aeb282 --- /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
