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

[Nick Sarnie via cfe-commits]

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

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

Signed-off-by: Sarnie, Nick <nick.sar...@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;
 }

>From 303fdae0e53dd1eb718591b002cc8d7347920734 Mon Sep 17 00:00:00 2001
From: "Sarnie, Nick" <nick.sar...@intel.com>
Date: Tue, 2 Sep 2025 14:01:28 -0700
Subject: [PATCH 2/2] fix comment

Signed-off-by: Sarnie, Nick <nick.sar...@intel.com>
---
 llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp 
b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index c9260e92513c9..7da2927be00f2 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -963,7 +963,7 @@ Constant *OpenMPIRBuilder::getOrCreateSrcLocStr(StringRef 
LocStr,
         return SrcLocStr = ConstantExpr::getPointerCast(&GV, Int8Ptr);
 
     SrcLocStr = Builder.CreateGlobalString(
-        LocStr, /* Name */ "",
+        LocStr, /*Name=*/ "",
         M.getDataLayout().getDefaultGlobalsAddressSpace(), &M);
   }
   return SrcLocStr;

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits