[llvm] e6c63d9 - [OMPIRBuilder] Use target global AS for SrcLocStr (#156520)

[via llvm-commits]

Author: Nick Sarnie
Date: 2025-09-03T14:35:33Z
New Revision: e6c63d920dec3e8874ac1dc3c3f19fb822f0ab06

URL: https://github.com/llvm/llvm-project/commit/e6c63d920dec3e8874ac1dc3c3f19fb822f0ab06
DIFF: https://github.com/llvm/llvm-project/commit/e6c63d920dec3e8874ac1dc3c3f19fb822f0ab06.diff

LOG: [OMPIRBuilder] Use target global AS for SrcLocStr (#156520)

We should set the correct target-specific AS for the SrcLocStr global
created in OMPIRBuilder.

We also may have to insert a constexpr addrspacecast because the struct
field type may be different than the value used to initialize it.

I actually want the cast to be from AS 1 to AS 4, but getting the type
to be AS4 relies on a PR currently in-review, so leave the cast target
to AS 0 for now.

---------

Signed-off-by: Sarnie, Nick <nick.sarnie at intel.com>

Added: 
    clang/test/OpenMP/spirv_locstr.cpp

Modified: 
    llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

Removed: 
    


################################################################################
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 03ea58318d4a9..3d5e487c8990f 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;
 }