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.sar...@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; } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
