[clang][OMPIRBuilder] Fix reduction codegen for SPIR-V #162529
+29
−9
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
llvmbot
added
clang
|
@llvm/pr-subscribers-clang-codegen Author: Nick Sarnie (sarnex) ChangesWhen creating function pointers, make sure the pointer is in the program address space. Also fix a spot where I forgot to set Full diff: https://github.com/llvm/llvm-project/pull/162529.diff 4 Files Affected:
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
index 4272d8b1a1f51..3613b6a143d42 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -869,6 +869,8 @@ CGOpenMPRuntimeGPU::CGOpenMPRuntimeGPU(CodeGenModule &CGM)
CGM.getLangOpts().OpenMPOffloadMandatory,
/*HasRequiresReverseOffload*/ false, /*HasRequiresUnifiedAddress*/ false,
hasRequiresUnifiedSharedMemory(), /*HasRequiresDynamicAllocators*/ false);
+ Config.setDefaultTargetAS(
+ CGM.getContext().getTargetInfo().getTargetAddressSpace(LangAS::Default));
OMPBuilder.setConfig(Config);
if (!CGM.getLangOpts().OpenMPIsTargetDevice)
@@ -1243,7 +1245,10 @@ void CGOpenMPRuntimeGPU::emitParallelCall(
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::Type *FnPtrTy = llvm::PointerType::get(
+ CGF.getLLVMContext(), CGM.getDataLayout().getProgramAddressSpace());
+
+ llvm::Value *FnPtr = Bld.CreateBitOrPointerCast(OutlinedFn, FnPtrTy);
// Create a private scope that will globalize the arguments
// passed from the outside of the target region.
diff --git a/clang/test/OpenMP/spirv_reduction.cpp b/clang/test/OpenMP/spirv_reduction.cpp
new file mode 100644
index 0000000000000..e0e7549de716c
--- /dev/null
+++ b/clang/test/OpenMP/spirv_reduction.cpp
@@ -0,0 +1,22 @@
+// 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: call spir_func addrspace(9) void @__kmpc_parallel_51(ptr addrspace(4) addrspacecast (ptr addrspace(1) @{{.*}} to ptr addrspace(4)),
+// CHECK-SAME: i32 %{{.*}}, i32 {{.*}}, i32 {{.*}}, i32 {{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(4) {{.*}}, ptr addrspace(4) %{{.*}}, i64 {{.*}})
+
+// CHECK: call addrspace(9) i32 @__kmpc_nvptx_teams_reduce_nowait_v2(ptr addrspace(4) addrspacecast (ptr addrspace(1) @{{.*}} to ptr addrspace(4)),
+// CHECK-SAME: ptr addrspace(4) %{{.*}}, i32 1024, i64 4, ptr addrspace(4) %{{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(9) @{{.*}})
+
+int main() {
+ int matrix_sum = 0;
+ #pragma omp target teams distribute parallel for \
+ reduction(+:matrix_sum) \
+ map(tofrom:matrix_sum)
+ for (int i = 0; i < 100; i++) {
+
+ }
+
+ return 0;
+}
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
index 01ca8da759ef7..1694a33510d79 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -42,6 +42,7 @@ __OMP_TYPE(Double)
OMP_TYPE(SizeTy, M.getDataLayout().getIntPtrType(Ctx))
OMP_TYPE(Int63, Type::getIntNTy(Ctx, 63))
+OMP_TYPE(FuncPtrTy, PointerType::get(Ctx, M.getDataLayout().getProgramAddressSpace()))
__OMP_PTR_TYPE(VoidPtr)
__OMP_PTR_TYPE(VoidPtrPtr)
@@ -471,7 +472,7 @@ __OMP_RTL(__kmpc_target_init, false, Int32, KernelEnvironmentPtr, KernelLaunchEn
__OMP_RTL(__kmpc_target_deinit, false, Void,)
__OMP_RTL(__kmpc_kernel_prepare_parallel, false, Void, VoidPtr)
__OMP_RTL(__kmpc_parallel_51, false, Void, IdentPtr, Int32, Int32, Int32, Int32,
- VoidPtr, VoidPtr, VoidPtrPtr, SizeTy)
+ FuncPtrTy, VoidPtr, VoidPtrPtr, SizeTy)
__OMP_RTL(__kmpc_for_static_loop_4, false, Void, IdentPtr, VoidPtr, VoidPtr, Int32, Int32, Int32, Int8)
__OMP_RTL(__kmpc_for_static_loop_4u, false, Void, IdentPtr, VoidPtr, VoidPtr, Int32, Int32, Int32, Int8)
__OMP_RTL(__kmpc_for_static_loop_8, false, Void, IdentPtr, VoidPtr, VoidPtr, Int64, Int64, Int64, Int8)
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 5980ee35a5cd2..286ed039b1214 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -3623,7 +3623,9 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createReductionsGPU(
// 1. Build a list of reduction variables.
// void *RedList[<n>] = {<ReductionVars>[0], ..., <ReductionVars>[<n>-1]};
auto Size = ReductionInfos.size();
- Type *PtrTy = PointerType::getUnqual(Ctx);
+ Type *PtrTy = PointerType::get(Ctx, Config.getDefaultTargetAS());
+ Type *FuncPtrTy =
+ Builder.getPtrTy(M.getDataLayout().getProgramAddressSpace());
Type *RedArrayTy = ArrayType::get(PtrTy, Size);
CodeGenIP = Builder.saveIP();
Builder.restoreIP(AllocaIP);
@@ -3667,9 +3669,9 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createReductionsGPU(
Builder.getInt64(MaxDataSize * ReductionInfos.size());
if (!IsTeamsReduction) {
Value *SarFuncCast =
- Builder.CreatePointerBitCastOrAddrSpaceCast(SarFunc, PtrTy);
+ Builder.CreatePointerBitCastOrAddrSpaceCast(SarFunc, FuncPtrTy);
Value *WcFuncCast =
- Builder.CreatePointerBitCastOrAddrSpaceCast(WcFunc, PtrTy);
+ Builder.CreatePointerBitCastOrAddrSpaceCast(WcFunc, FuncPtrTy);
Value *Args[] = {SrcLocInfo, ReductionDataSize, RL, SarFuncCast,
WcFuncCast};
Function *Pv2Ptr = getOrCreateRuntimeFunctionPtr(
@@ -10072,13 +10074,14 @@ void OpenMPIRBuilder::initializeTypes(Module &M) {
LLVMContext &Ctx = M.getContext();
StructType *T;
unsigned DefaultTargetAS = Config.getDefaultTargetAS();
+ unsigned ProgramAS = M.getDataLayout().getProgramAddressSpace();
#define OMP_TYPE(VarName, InitValue) VarName = InitValue;
#define OMP_ARRAY_TYPE(VarName, ElemTy, ArraySize) \
VarName##Ty = ArrayType::get(ElemTy, ArraySize); \
VarName##PtrTy = PointerType::get(Ctx, DefaultTargetAS);
#define OMP_FUNCTION_TYPE(VarName, IsVarArg, ReturnType, ...) \
VarName = FunctionType::get(ReturnType, {__VA_ARGS__}, IsVarArg); \
- VarName##Ptr = PointerType::get(Ctx, DefaultTargetAS);
+ VarName##Ptr = PointerType::get(Ctx, ProgramAS);
#define OMP_STRUCT_TYPE(VarName, StructName, Packed, ...) \
T = StructType::getTypeByName(Ctx, StructName); \
if (!T) \
|
|
@llvm/pr-subscribers-flang-openmp Author: Nick Sarnie (sarnex) ChangesWhen creating function pointers, make sure the pointer is in the program address space. Also fix a spot where I forgot to set Full diff: https://github.com/llvm/llvm-project/pull/162529.diff 4 Files Affected:
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
index 4272d8b1a1f51..3613b6a143d42 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -869,6 +869,8 @@ CGOpenMPRuntimeGPU::CGOpenMPRuntimeGPU(CodeGenModule &CGM)
CGM.getLangOpts().OpenMPOffloadMandatory,
/*HasRequiresReverseOffload*/ false, /*HasRequiresUnifiedAddress*/ false,
hasRequiresUnifiedSharedMemory(), /*HasRequiresDynamicAllocators*/ false);
+ Config.setDefaultTargetAS(
+ CGM.getContext().getTargetInfo().getTargetAddressSpace(LangAS::Default));
OMPBuilder.setConfig(Config);
if (!CGM.getLangOpts().OpenMPIsTargetDevice)
@@ -1243,7 +1245,10 @@ void CGOpenMPRuntimeGPU::emitParallelCall(
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::Type *FnPtrTy = llvm::PointerType::get(
+ CGF.getLLVMContext(), CGM.getDataLayout().getProgramAddressSpace());
+
+ llvm::Value *FnPtr = Bld.CreateBitOrPointerCast(OutlinedFn, FnPtrTy);
// Create a private scope that will globalize the arguments
// passed from the outside of the target region.
diff --git a/clang/test/OpenMP/spirv_reduction.cpp b/clang/test/OpenMP/spirv_reduction.cpp
new file mode 100644
index 0000000000000..e0e7549de716c
--- /dev/null
+++ b/clang/test/OpenMP/spirv_reduction.cpp
@@ -0,0 +1,22 @@
+// 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: call spir_func addrspace(9) void @__kmpc_parallel_51(ptr addrspace(4) addrspacecast (ptr addrspace(1) @{{.*}} to ptr addrspace(4)),
+// CHECK-SAME: i32 %{{.*}}, i32 {{.*}}, i32 {{.*}}, i32 {{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(4) {{.*}}, ptr addrspace(4) %{{.*}}, i64 {{.*}})
+
+// CHECK: call addrspace(9) i32 @__kmpc_nvptx_teams_reduce_nowait_v2(ptr addrspace(4) addrspacecast (ptr addrspace(1) @{{.*}} to ptr addrspace(4)),
+// CHECK-SAME: ptr addrspace(4) %{{.*}}, i32 1024, i64 4, ptr addrspace(4) %{{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(9) @{{.*}})
+
+int main() {
+ int matrix_sum = 0;
+ #pragma omp target teams distribute parallel for \
+ reduction(+:matrix_sum) \
+ map(tofrom:matrix_sum)
+ for (int i = 0; i < 100; i++) {
+
+ }
+
+ return 0;
+}
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
index 01ca8da759ef7..1694a33510d79 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -42,6 +42,7 @@ __OMP_TYPE(Double)
OMP_TYPE(SizeTy, M.getDataLayout().getIntPtrType(Ctx))
OMP_TYPE(Int63, Type::getIntNTy(Ctx, 63))
+OMP_TYPE(FuncPtrTy, PointerType::get(Ctx, M.getDataLayout().getProgramAddressSpace()))
__OMP_PTR_TYPE(VoidPtr)
__OMP_PTR_TYPE(VoidPtrPtr)
@@ -471,7 +472,7 @@ __OMP_RTL(__kmpc_target_init, false, Int32, KernelEnvironmentPtr, KernelLaunchEn
__OMP_RTL(__kmpc_target_deinit, false, Void,)
__OMP_RTL(__kmpc_kernel_prepare_parallel, false, Void, VoidPtr)
__OMP_RTL(__kmpc_parallel_51, false, Void, IdentPtr, Int32, Int32, Int32, Int32,
- VoidPtr, VoidPtr, VoidPtrPtr, SizeTy)
+ FuncPtrTy, VoidPtr, VoidPtrPtr, SizeTy)
__OMP_RTL(__kmpc_for_static_loop_4, false, Void, IdentPtr, VoidPtr, VoidPtr, Int32, Int32, Int32, Int8)
__OMP_RTL(__kmpc_for_static_loop_4u, false, Void, IdentPtr, VoidPtr, VoidPtr, Int32, Int32, Int32, Int8)
__OMP_RTL(__kmpc_for_static_loop_8, false, Void, IdentPtr, VoidPtr, VoidPtr, Int64, Int64, Int64, Int8)
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 5980ee35a5cd2..286ed039b1214 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -3623,7 +3623,9 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createReductionsGPU(
// 1. Build a list of reduction variables.
// void *RedList[<n>] = {<ReductionVars>[0], ..., <ReductionVars>[<n>-1]};
auto Size = ReductionInfos.size();
- Type *PtrTy = PointerType::getUnqual(Ctx);
+ Type *PtrTy = PointerType::get(Ctx, Config.getDefaultTargetAS());
+ Type *FuncPtrTy =
+ Builder.getPtrTy(M.getDataLayout().getProgramAddressSpace());
Type *RedArrayTy = ArrayType::get(PtrTy, Size);
CodeGenIP = Builder.saveIP();
Builder.restoreIP(AllocaIP);
@@ -3667,9 +3669,9 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createReductionsGPU(
Builder.getInt64(MaxDataSize * ReductionInfos.size());
if (!IsTeamsReduction) {
Value *SarFuncCast =
- Builder.CreatePointerBitCastOrAddrSpaceCast(SarFunc, PtrTy);
+ Builder.CreatePointerBitCastOrAddrSpaceCast(SarFunc, FuncPtrTy);
Value *WcFuncCast =
- Builder.CreatePointerBitCastOrAddrSpaceCast(WcFunc, PtrTy);
+ Builder.CreatePointerBitCastOrAddrSpaceCast(WcFunc, FuncPtrTy);
Value *Args[] = {SrcLocInfo, ReductionDataSize, RL, SarFuncCast,
WcFuncCast};
Function *Pv2Ptr = getOrCreateRuntimeFunctionPtr(
@@ -10072,13 +10074,14 @@ void OpenMPIRBuilder::initializeTypes(Module &M) {
LLVMContext &Ctx = M.getContext();
StructType *T;
unsigned DefaultTargetAS = Config.getDefaultTargetAS();
+ unsigned ProgramAS = M.getDataLayout().getProgramAddressSpace();
#define OMP_TYPE(VarName, InitValue) VarName = InitValue;
#define OMP_ARRAY_TYPE(VarName, ElemTy, ArraySize) \
VarName##Ty = ArrayType::get(ElemTy, ArraySize); \
VarName##PtrTy = PointerType::get(Ctx, DefaultTargetAS);
#define OMP_FUNCTION_TYPE(VarName, IsVarArg, ReturnType, ...) \
VarName = FunctionType::get(ReturnType, {__VA_ARGS__}, IsVarArg); \
- VarName##Ptr = PointerType::get(Ctx, DefaultTargetAS);
+ VarName##Ptr = PointerType::get(Ctx, ProgramAS);
#define OMP_STRUCT_TYPE(VarName, StructName, Packed, ...) \
T = StructType::getTypeByName(Ctx, StructName); \
if (!T) \
|
|
@llvm/pr-subscribers-clang Author: Nick Sarnie (sarnex) ChangesWhen creating function pointers, make sure the pointer is in the program address space. Also fix a spot where I forgot to set Full diff: https://github.com/llvm/llvm-project/pull/162529.diff 4 Files Affected:
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
index 4272d8b1a1f51..3613b6a143d42 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -869,6 +869,8 @@ CGOpenMPRuntimeGPU::CGOpenMPRuntimeGPU(CodeGenModule &CGM)
CGM.getLangOpts().OpenMPOffloadMandatory,
/*HasRequiresReverseOffload*/ false, /*HasRequiresUnifiedAddress*/ false,
hasRequiresUnifiedSharedMemory(), /*HasRequiresDynamicAllocators*/ false);
+ Config.setDefaultTargetAS(
+ CGM.getContext().getTargetInfo().getTargetAddressSpace(LangAS::Default));
OMPBuilder.setConfig(Config);
if (!CGM.getLangOpts().OpenMPIsTargetDevice)
@@ -1243,7 +1245,10 @@ void CGOpenMPRuntimeGPU::emitParallelCall(
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::Type *FnPtrTy = llvm::PointerType::get(
+ CGF.getLLVMContext(), CGM.getDataLayout().getProgramAddressSpace());
+
+ llvm::Value *FnPtr = Bld.CreateBitOrPointerCast(OutlinedFn, FnPtrTy);
// Create a private scope that will globalize the arguments
// passed from the outside of the target region.
diff --git a/clang/test/OpenMP/spirv_reduction.cpp b/clang/test/OpenMP/spirv_reduction.cpp
new file mode 100644
index 0000000000000..e0e7549de716c
--- /dev/null
+++ b/clang/test/OpenMP/spirv_reduction.cpp
@@ -0,0 +1,22 @@
+// 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: call spir_func addrspace(9) void @__kmpc_parallel_51(ptr addrspace(4) addrspacecast (ptr addrspace(1) @{{.*}} to ptr addrspace(4)),
+// CHECK-SAME: i32 %{{.*}}, i32 {{.*}}, i32 {{.*}}, i32 {{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(4) {{.*}}, ptr addrspace(4) %{{.*}}, i64 {{.*}})
+
+// CHECK: call addrspace(9) i32 @__kmpc_nvptx_teams_reduce_nowait_v2(ptr addrspace(4) addrspacecast (ptr addrspace(1) @{{.*}} to ptr addrspace(4)),
+// CHECK-SAME: ptr addrspace(4) %{{.*}}, i32 1024, i64 4, ptr addrspace(4) %{{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(9) @{{.*}}, ptr addrspace(9) @{{.*}})
+
+int main() {
+ int matrix_sum = 0;
+ #pragma omp target teams distribute parallel for \
+ reduction(+:matrix_sum) \
+ map(tofrom:matrix_sum)
+ for (int i = 0; i < 100; i++) {
+
+ }
+
+ return 0;
+}
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
index 01ca8da759ef7..1694a33510d79 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -42,6 +42,7 @@ __OMP_TYPE(Double)
OMP_TYPE(SizeTy, M.getDataLayout().getIntPtrType(Ctx))
OMP_TYPE(Int63, Type::getIntNTy(Ctx, 63))
+OMP_TYPE(FuncPtrTy, PointerType::get(Ctx, M.getDataLayout().getProgramAddressSpace()))
__OMP_PTR_TYPE(VoidPtr)
__OMP_PTR_TYPE(VoidPtrPtr)
@@ -471,7 +472,7 @@ __OMP_RTL(__kmpc_target_init, false, Int32, KernelEnvironmentPtr, KernelLaunchEn
__OMP_RTL(__kmpc_target_deinit, false, Void,)
__OMP_RTL(__kmpc_kernel_prepare_parallel, false, Void, VoidPtr)
__OMP_RTL(__kmpc_parallel_51, false, Void, IdentPtr, Int32, Int32, Int32, Int32,
- VoidPtr, VoidPtr, VoidPtrPtr, SizeTy)
+ FuncPtrTy, VoidPtr, VoidPtrPtr, SizeTy)
__OMP_RTL(__kmpc_for_static_loop_4, false, Void, IdentPtr, VoidPtr, VoidPtr, Int32, Int32, Int32, Int8)
__OMP_RTL(__kmpc_for_static_loop_4u, false, Void, IdentPtr, VoidPtr, VoidPtr, Int32, Int32, Int32, Int8)
__OMP_RTL(__kmpc_for_static_loop_8, false, Void, IdentPtr, VoidPtr, VoidPtr, Int64, Int64, Int64, Int8)
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 5980ee35a5cd2..286ed039b1214 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -3623,7 +3623,9 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createReductionsGPU(
// 1. Build a list of reduction variables.
// void *RedList[<n>] = {<ReductionVars>[0], ..., <ReductionVars>[<n>-1]};
auto Size = ReductionInfos.size();
- Type *PtrTy = PointerType::getUnqual(Ctx);
+ Type *PtrTy = PointerType::get(Ctx, Config.getDefaultTargetAS());
+ Type *FuncPtrTy =
+ Builder.getPtrTy(M.getDataLayout().getProgramAddressSpace());
Type *RedArrayTy = ArrayType::get(PtrTy, Size);
CodeGenIP = Builder.saveIP();
Builder.restoreIP(AllocaIP);
@@ -3667,9 +3669,9 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createReductionsGPU(
Builder.getInt64(MaxDataSize * ReductionInfos.size());
if (!IsTeamsReduction) {
Value *SarFuncCast =
- Builder.CreatePointerBitCastOrAddrSpaceCast(SarFunc, PtrTy);
+ Builder.CreatePointerBitCastOrAddrSpaceCast(SarFunc, FuncPtrTy);
Value *WcFuncCast =
- Builder.CreatePointerBitCastOrAddrSpaceCast(WcFunc, PtrTy);
+ Builder.CreatePointerBitCastOrAddrSpaceCast(WcFunc, FuncPtrTy);
Value *Args[] = {SrcLocInfo, ReductionDataSize, RL, SarFuncCast,
WcFuncCast};
Function *Pv2Ptr = getOrCreateRuntimeFunctionPtr(
@@ -10072,13 +10074,14 @@ void OpenMPIRBuilder::initializeTypes(Module &M) {
LLVMContext &Ctx = M.getContext();
StructType *T;
unsigned DefaultTargetAS = Config.getDefaultTargetAS();
+ unsigned ProgramAS = M.getDataLayout().getProgramAddressSpace();
#define OMP_TYPE(VarName, InitValue) VarName = InitValue;
#define OMP_ARRAY_TYPE(VarName, ElemTy, ArraySize) \
VarName##Ty = ArrayType::get(ElemTy, ArraySize); \
VarName##PtrTy = PointerType::get(Ctx, DefaultTargetAS);
#define OMP_FUNCTION_TYPE(VarName, IsVarArg, ReturnType, ...) \
VarName = FunctionType::get(ReturnType, {__VA_ARGS__}, IsVarArg); \
- VarName##Ptr = PointerType::get(Ctx, DefaultTargetAS);
+ VarName##Ptr = PointerType::get(Ctx, ProgramAS);
#define OMP_STRUCT_TYPE(VarName, StructName, Packed, ...) \
T = StructType::getTypeByName(Ctx, StructName); \
if (!T) \
|
jhuber6
approved these changes
Oct 8, 2025
| @@ -0,0 +1,22 @@ | |||
| // 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 | |||
There was a problem hiding this comment.
Is there really not another reduction test we can just slap spirv onto?
There was a problem hiding this comment.
Seems reasonable, I'm assuming the other targets have the same behavior as before since we don't split on function AS.
Right, for other targets default target AS and program AS are 0 so none of these addrspace changes have any effect
Is there really not another reduction test we can just slap spirv onto?
Good point, I'm sure there is, let me take a look
There was a problem hiding this comment.
Just pushed a commit that modifies an existing test instead, thanks
Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
When creating function pointers, make sure the pointer is in the program address space.
Also fix a spot where I forgot to set
setDefaultTargetASand one spot where we didn't use the default AS for a pointer.