[clang] [llvm] [OMPIRBuilder] Use target global AS for SrcLocStr (PR #156520)

[Nick Sarnie via llvm-commits]

https://github.com/sarnex created https://github.com/llvm/llvm-project/pull/156520

None

>From 9c502a4fad10d8854d621ff72f08df9bd31f8949 Mon Sep 17 00:00:00 2001
From: "Sarnie, Nick" <nick.sarnie at intel.com>
Date: Tue, 2 Sep 2025 12:45:22 -0700
Subject: [PATCH] [clang][OpenMP] Use target global AS for SrcLocStr

Signed-off-by: Sarnie, Nick <nick.sarnie at intel.com>
---
 clang/test/OpenMP/spirv_locstr.cpp        | 15 +++++++++++++++
 llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 12 ++++++++++--
 2 files changed, 25 insertions(+), 2 deletions(-)
 create mode 100644 clang/test/OpenMP/spirv_locstr.cpp

diff --git a/clang/test/OpenMP/spirv_locstr.cpp b/clang/test/OpenMP/spirv_locstr.cpp
new file mode 100644
index 0000000000000..20d9c9d2f7393
--- /dev/null
+++ b/clang/test/OpenMP/spirv_locstr.cpp
@@ -0,0 +1,15 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=spirv64-intel -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple spirv64-intel -fopenmp-targets=spirv64-intel -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck %s
+
+// expected-no-diagnostics
+
+// CHECK: @[[#LOC:]] = private unnamed_addr addrspace(1) constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1
+// CHECK: = private unnamed_addr addrspace(1) constant %struct.ident_t { i32 0, i32 2, i32 0, i32 {{.*}}, ptr addrspacecast (ptr addrspace(1) @[[#LOC]] to ptr) }, align 8
+
+int main() {
+  int ret = 0;
+  #pragma omp target
+  for(int i = 0; i < 5; i++)
+    ret++;
+  return ret;
+}
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index e740c2819fec9..c9260e92513c9 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -915,6 +915,13 @@ Constant *OpenMPIRBuilder::getOrCreateIdent(Constant *SrcLocStr,
                              ConstantInt::get(Int32, uint32_t(LocFlags)),
                              ConstantInt::get(Int32, Reserve2Flags),
                              ConstantInt::get(Int32, SrcLocStrSize), SrcLocStr};
+
+    size_t SrcLocStrArgIdx = 4;
+    if (OpenMPIRBuilder::Ident->getElementType(SrcLocStrArgIdx)
+            ->getPointerAddressSpace() !=
+        IdentData[SrcLocStrArgIdx]->getType()->getPointerAddressSpace())
+      IdentData[SrcLocStrArgIdx] = ConstantExpr::getAddrSpaceCast(
+          SrcLocStr, OpenMPIRBuilder::Ident->getElementType(SrcLocStrArgIdx));
     Constant *Initializer =
         ConstantStruct::get(OpenMPIRBuilder::Ident, IdentData);
 
@@ -955,8 +962,9 @@ Constant *OpenMPIRBuilder::getOrCreateSrcLocStr(StringRef LocStr,
           GV.getInitializer() == Initializer)
         return SrcLocStr = ConstantExpr::getPointerCast(&GV, Int8Ptr);
 
-    SrcLocStr = Builder.CreateGlobalString(LocStr, /* Name */ "",
-                                           /* AddressSpace */ 0, &M);
+    SrcLocStr = Builder.CreateGlobalString(
+        LocStr, /* Name */ "",
+        M.getDataLayout().getDefaultGlobalsAddressSpace(), &M);
   }
   return SrcLocStr;
 }