https://github.com/sarnex updated
https://github.com/llvm/llvm-project/pull/135251
>From 675dde092c16a779d858f6082d0aab19acae3a1e Mon Sep 17 00:00:00 2001
From: "Sarnie, Nick" <nick.sar...@intel.com>
Date: Wed, 9 Apr 2025 08:42:49 -0700
Subject: [PATCH 1/2] [clang][OpenMP][SPIR-V] Fix addrspace of globals
Signed-off-by: Sarnie, Nick <nick.sar...@intel.com>
---
clang/lib/Basic/Targets/SPIR.h | 10 ++++---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 28 +++++++++++++++----
clang/lib/CodeGen/CGOpenMPRuntime.h | 5 ++++
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp | 21 ++++++++------
clang/lib/CodeGen/CGStmtOpenMP.cpp | 2 ++
clang/test/OpenMP/spirv_target_addrspace.c | 21 ++++++++++++++
.../test/OpenMP/spirv_target_addrspace_simd.c | 23 +++++++++++++++
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 13 +++++++--
8 files changed, 102 insertions(+), 21 deletions(-)
create mode 100644 clang/test/OpenMP/spirv_target_addrspace.c
create mode 100644 clang/test/OpenMP/spirv_target_addrspace_simd.c
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;
}
>From e36aa2c3de20b6aab01463a782f9e1b7c9883a89 Mon Sep 17 00:00:00 2001
From: "Sarnie, Nick" <nick.sar...@intel.com>
Date: Mon, 21 Apr 2025 11:14:26 -0700
Subject: [PATCH 2/2] add cuda spirv test
Signed-off-by: Sarnie, Nick <nick.sar...@intel.com>
---
clang/test/CodeGenCUDASPIRV/printf.cu | 11 +++++++++++
1 file changed, 11 insertions(+)
create mode 100644 clang/test/CodeGenCUDASPIRV/printf.cu
diff --git a/clang/test/CodeGenCUDASPIRV/printf.cu
b/clang/test/CodeGenCUDASPIRV/printf.cu
new file mode 100644
index 0000000000000..936e920f4a755
--- /dev/null
+++ b/clang/test/CodeGenCUDASPIRV/printf.cu
@@ -0,0 +1,11 @@
+// RUN: %clang_cc1 -fcuda-is-device -triple spirv32 -o - -emit-llvm -x cuda %s
| FileCheck --check-prefix=CHECK-SPIRV32 %s
+// RUN: %clang_cc1 -fcuda-is-device -triple spirv64 -o - -emit-llvm -x cuda %s
| FileCheck --check-prefix=CHECK-SPIRV64 %s
+
+// CHECK-SPIRV32: @.str = private unnamed_addr addrspace(4) constant [13 x i8]
c"Hello World\0A\00", align 1
+// CHECK-SPIRV64: @.str = private unnamed_addr addrspace(1) constant [13 x i8]
c"Hello World\0A\00", align 1
+
+extern "C" __attribute__((device)) int printf(const char* format, ...);
+
+__attribute__((global)) void printf_kernel() {
+ printf("Hello World\n");
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
