[clang] [llvm] [clang][OpenMP][SPIR-V] Fix AS of globals and set the default AS to 4 (PR #135251)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Apr 11 09:47:59 PDT 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Nick Sarnie (sarnex)
<details>
<summary>Changes</summary>
Based on feedback from https://github.com/llvm/llvm-project/pull/134399, we use the address space map that sets the default AS to 4 for OpenMP SPIR-V offload.
The new AS map had the wrong target AS for `opencl_global` and `opencl_constant`, which is what we end up using to get the AS during OpenMP target codegen, so I updated it to match the old default AS 0 map.
After this PR, I will work on simplifying the condition of the old default AS 0 map to eventually only be OCL with no generic addrspace, but there are many failures so I wanted to do it step by step, and this is the first one.
There are relatively minor changes to OpenMP codegen, mostly just addrspacecasts (because globals are AS 1 in SPIR-V, so we need to cast to AS 4/no AS somewhat often) or use the correct address space to create a global.
---
Full diff: https://github.com/llvm/llvm-project/pull/135251.diff
8 Files Affected:
- (modified) clang/lib/Basic/Targets/SPIR.h (+6-4)
- (modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+22-6)
- (modified) clang/lib/CodeGen/CGOpenMPRuntime.h (+5)
- (modified) clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp (+12-9)
- (modified) clang/lib/CodeGen/CGStmtOpenMP.cpp (+2)
- (added) clang/test/OpenMP/spirv_target_addrspace.c (+21)
- (added) clang/test/OpenMP/spirv_target_addrspace_simd.c (+23)
- (modified) llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp (+11-2)
``````````diff
diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h
index 5ea727364d24b..0f4f74ac95749 100644
--- a/clang/lib/Basic/Targets/SPIR.h
+++ b/clang/lib/Basic/Targets/SPIR.h
@@ -57,10 +57,11 @@ static const unsigned SPIRDefIsPrivMap[] = {
// Used by both the SPIR and SPIR-V targets.
static const unsigned SPIRDefIsGenMap[] = {
4, // Default
- // OpenCL address space values for this map are dummy and they can't be used
- 0, // opencl_global
+ // Some OpenCL address space values for this map are dummy and they can't be
+ // used
+ 1, // opencl_global
0, // opencl_local
- 0, // opencl_constant
+ 2, // opencl_constant
0, // opencl_private
0, // opencl_generic
0, // opencl_global_device
@@ -216,7 +217,8 @@ class LLVM_LIBRARY_VISIBILITY BaseSPIRTargetInfo : public TargetInfo {
/*DefaultIsGeneric=*/Opts.SYCLIsDevice ||
// The address mapping from HIP/CUDA language for device code is only
// defined for SPIR-V.
- (getTriple().isSPIRV() && Opts.CUDAIsDevice));
+ (getTriple().isSPIRV() &&
+ (Opts.CUDAIsDevice || Opts.OpenMPIsTargetDevice)));
}
void setSupportedOpenCLOpts() override {
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 5736864d4cc6b..5780f1ded3259 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -2524,6 +2524,16 @@ void CGOpenMPRuntime::emitForDispatchInit(
Args);
}
+llvm::Value *CGOpenMPRuntime::createRuntimeFunctionArgAddrSpaceCast(
+ CodeGenFunction &CGF, llvm::FunctionCallee RuntimeFcn, size_t ArgIdx,
+ llvm::Value *Arg) {
+ llvm::Type *ParamTy = RuntimeFcn.getFunctionType()->getParamType(ArgIdx);
+ llvm::Type *ArgTy = Arg->getType();
+ if (!ParamTy->isPointerTy())
+ return Arg;
+ return CGF.Builder.CreateAddrSpaceCast(Arg, ParamTy);
+}
+
void CGOpenMPRuntime::emitForDispatchDeinit(CodeGenFunction &CGF,
SourceLocation Loc) {
if (!CGF.HaveInsertPoint())
@@ -2572,12 +2582,18 @@ static void emitForStaticInitCall(
ThreadId,
CGF.Builder.getInt32(addMonoNonMonoModifier(CGF.CGM, Schedule, M1,
M2)), // Schedule type
- Values.IL.emitRawPointer(CGF), // &isLastIter
- Values.LB.emitRawPointer(CGF), // &LB
- Values.UB.emitRawPointer(CGF), // &UB
- Values.ST.emitRawPointer(CGF), // &Stride
- CGF.Builder.getIntN(Values.IVSize, 1), // Incr
- Chunk // Chunk
+ CGOpenMPRuntime::createRuntimeFunctionArgAddrSpaceCast(
+ CGF, ForStaticInitFunction, 3,
+ Values.IL.emitRawPointer(CGF)), // &isLastIter
+ CGOpenMPRuntime::createRuntimeFunctionArgAddrSpaceCast(
+ CGF, ForStaticInitFunction, 4, Values.LB.emitRawPointer(CGF)), // &LB
+ CGOpenMPRuntime::createRuntimeFunctionArgAddrSpaceCast(
+ CGF, ForStaticInitFunction, 5, Values.UB.emitRawPointer(CGF)), // &UB
+ CGOpenMPRuntime::createRuntimeFunctionArgAddrSpaceCast(
+ CGF, ForStaticInitFunction, 6,
+ Values.ST.emitRawPointer(CGF)), // &Stride
+ CGF.Builder.getIntN(Values.IVSize, 1), // Incr
+ Chunk // Chunk
};
CGF.EmitRuntimeCall(ForStaticInitFunction, Args);
}
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.h b/clang/lib/CodeGen/CGOpenMPRuntime.h
index 4321712e1521d..c918c77b4266c 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.h
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.h
@@ -1657,6 +1657,11 @@ class CGOpenMPRuntime {
/// Returns true if the variable is a local variable in untied task.
bool isLocalVarInUntiedTask(CodeGenFunction &CGF, const VarDecl *VD) const;
+
+ static llvm::Value *
+ createRuntimeFunctionArgAddrSpaceCast(CodeGenFunction &CGF,
+ llvm::FunctionCallee RuntimeFcn,
+ size_t ArgIdx, llvm::Value *Arg);
};
/// Class supports emissionof SIMD-only code.
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
index f697c13f4c522..0bfa49dee0c53 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -1217,11 +1217,13 @@ void CGOpenMPRuntimeGPU::emitParallelCall(CodeGenFunction &CGF,
CGBuilderTy &Bld = CGF.Builder;
llvm::Value *NumThreadsVal = NumThreads;
llvm::Function *WFn = WrapperFunctionsMap[OutlinedFn];
+ llvm::FunctionCallee RuntimeFn = OMPBuilder.getOrCreateRuntimeFunction(
+ CGM.getModule(), OMPRTL___kmpc_parallel_51);
llvm::Value *ID = llvm::ConstantPointerNull::get(CGM.Int8PtrTy);
if (WFn)
ID = Bld.CreateBitOrPointerCast(WFn, CGM.Int8PtrTy);
- llvm::Value *FnPtr = Bld.CreateBitOrPointerCast(OutlinedFn, CGM.Int8PtrTy);
-
+ llvm::Value *FnPtr = Bld.CreateAddrSpaceCast(OutlinedFn, CGM.Int8PtrTy);
+ FnPtr = Bld.CreateBitOrPointerCast(FnPtr, CGM.Int8PtrTy);
// Create a private scope that will globalize the arguments
// passed from the outside of the target region.
// TODO: Is that needed?
@@ -1268,14 +1270,15 @@ void CGOpenMPRuntimeGPU::emitParallelCall(CodeGenFunction &CGF,
IfCondVal,
NumThreadsVal,
llvm::ConstantInt::get(CGF.Int32Ty, -1),
- FnPtr,
- ID,
- Bld.CreateBitOrPointerCast(CapturedVarsAddrs.emitRawPointer(CGF),
- CGF.VoidPtrPtrTy),
+ createRuntimeFunctionArgAddrSpaceCast(CGF, RuntimeFn, 5, FnPtr),
+ createRuntimeFunctionArgAddrSpaceCast(CGF, RuntimeFn, 6, ID),
+ createRuntimeFunctionArgAddrSpaceCast(
+ CGF, RuntimeFn, 7,
+ Bld.CreateBitOrPointerCast(CapturedVarsAddrs.emitRawPointer(CGF),
+ CGF.VoidPtrPtrTy)),
llvm::ConstantInt::get(CGM.SizeTy, CapturedVars.size())};
- CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
- CGM.getModule(), OMPRTL___kmpc_parallel_51),
- Args);
+
+ CGF.EmitRuntimeCall(RuntimeFn, Args);
};
RegionCodeGenTy RCG(ParallelGen);
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index 156f64bb5f508..78fd65750fc02 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -585,6 +585,8 @@ static llvm::Function *emitOutlinedFunctionPrologue(
F->removeFnAttr(llvm::Attribute::NoInline);
F->addFnAttr(llvm::Attribute::AlwaysInline);
}
+ if (CGM.getTriple().isSPIRV())
+ F->setCallingConv(llvm::CallingConv::SPIR_FUNC);
// Generate the function.
CGF.StartFunction(CD, Ctx.VoidTy, F, FuncInfo, TargetArgs,
diff --git a/clang/test/OpenMP/spirv_target_addrspace.c b/clang/test/OpenMP/spirv_target_addrspace.c
new file mode 100644
index 0000000000000..9e5eeff73eed6
--- /dev/null
+++ b/clang/test/OpenMP/spirv_target_addrspace.c
@@ -0,0 +1,21 @@
+// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=spirv64 -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=spirv64 -fopenmp-is-target-device -triple spirv64 -fopenmp-host-ir-file-path %t-host.bc -emit-llvm %s -o - | FileCheck %s
+
+extern int fcn(const char[]);
+
+#pragma omp declare target
+// CHECK: @global = addrspace(1) global i32 0, align 4
+// CHECK: @.str = private unnamed_addr addrspace(1) constant [4 x i8] c"foo\00", align 1
+int global = 0;
+#pragma omp end declare target
+int main() {
+ // CHECK: = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @__omp_offloading_{{.*}}_kernel_environment to ptr), ptr %{{.*}})
+ #pragma omp target
+ {
+ for(int i = 0; i < 1024; i++)
+ global++;
+ fcn("foo");
+ }
+ return global;
+}
+
diff --git a/clang/test/OpenMP/spirv_target_addrspace_simd.c b/clang/test/OpenMP/spirv_target_addrspace_simd.c
new file mode 100644
index 0000000000000..31b00ab555596
--- /dev/null
+++ b/clang/test/OpenMP/spirv_target_addrspace_simd.c
@@ -0,0 +1,23 @@
+// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=spirv64 -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=spirv64 -fopenmp-is-target-device -triple spirv64 -fopenmp-host-ir-file-path %t-host.bc -emit-llvm %s -o - | FileCheck %s
+
+int main() {
+ int x = 0;
+
+#pragma omp target teams distribute parallel for simd
+ for(int i = 0; i < 1024; i++)
+ x+=i;
+ return x;
+}
+
+// CHECK: @[[#STRLOC:]] = private unnamed_addr addrspace(1) constant [{{.*}} x i8] c{{.*}}, align 1
+// CHECK: @[[#IDENT:]] = private unnamed_addr addrspace(1) constant %struct.ident_t { i32 {{.*}}, i32 2050, i32 {{.*}}, i32 {{.*}}, ptr addrspacecast (ptr addrspace(1) @[[#STRLOC]] to ptr) }, align 8
+// CHECK: define internal spir_func void @__omp_offloading_{{.*}}_omp_outlined(ptr addrspace(4) noalias noundef {{.*}}., ptr addrspace(4) noalias noundef {{.*}}, i64 noundef {{.*}}) #{{.*}} {
+// CHECK: = load ptr addrspace(4), ptr addrspace(4) %{{.*}}, align 8
+// CHECK: = load i32, ptr addrspace(4) %{{.*}}, align 4
+// CHECK: = addrspacecast ptr addrspace(4) %{{.*}} to ptr
+// CHECK: = addrspacecast ptr addrspace(4) %{{.*}} to ptr
+// CHECK: = addrspacecast ptr addrspace(4) %{{.*}} to ptr
+// CHECK: = addrspacecast ptr addrspace(4) %{{.*}} to ptr
+// CHECK: call spir_func void @__kmpc_distribute_static_init{{.*}}(ptr addrspacecast (ptr addrspace(1) @[[#IDENT]] to ptr), i32 %{{.*}}, i32 {{.*}}, ptr %{{.*}}, ptr %{{.*}}, ptr %{{.*}}, ptr %{{.*}}, i32 {{.*}}, i32 %{{.*}})
+// CHECK: call spir_func void @__kmpc_distribute_static_fini{{.*}}(ptr addrspacecast (ptr addrspace(1) @[[#IDENT]] to ptr), i32 %{{.*}})
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 13b727d226738..e7dc82acb9201 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -910,6 +910,14 @@ 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);
@@ -950,8 +958,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;
}
``````````
</details>
https://github.com/llvm/llvm-project/pull/135251
More information about the cfe-commits
mailing list