[clang] e34e311 - [clang][OpenMP] Cast by-reference captures to the address space of the captured field (#226229)

[via cfe-commits]

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