[clang][OpenMP][SPIR-V] Fix addrspace of pointer kernel arguments #157172
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
sarnex
commented
Sep 8, 2025
clang/lib/CodeGen/CGStmtOpenMP.cpp
Outdated
Member
Author
There was a problem hiding this comment.
It took me a while to find kernel detection logic that passes all the tests, this seems sane to be but I am definitely not an expert in this code, so would appreciate a closer look at this part, feel free to add anyone to the review, thanks!
Contributor
There was a problem hiding this comment.
It's tough to tell outside of the LLVM-IR because we generate many things to kernels. What's here seems to make sense to me, but I'm assuming that the check on target applies to all the variants of target.
Member
Author
There was a problem hiding this comment.
Yep, isOpenMPTargetExecutionDirective is a recursive check so any variant of a target statement should return true here.
bool clang::isOpenMPTargetExecutionDirective(OpenMPDirectiveKind DKind) {
return DKind == OMPD_target ||
llvm::is_contained(getLeafConstructs(DKind), OMPD_target);
}
Thanks for the review!
Signed-off-by: Sarnie, Nick <[email protected]>
llvmbot
added
clang
Member
|
@llvm/pr-subscribers-clang-codegen @llvm/pr-subscribers-clang Author: Nick Sarnie (sarnex) ChangesIn SPIR-V, kernel arguments are not allowed to be in the Generic AS, in both Intel's internal SPIR-V offloading implementation as well as HIPSPV, Currently with Generic AS the To get the existing logic to set the correct AS to kick in, we need to know if the function is a kernel or not at the time we first create the function that may end up as the kernel. I use the existing Rename the function to be accurate to the current behavior, Full diff: https://github.com/llvm/llvm-project/pull/157172.diff 8 Files Affected:
diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp
index a94a7ed51521c..0b2fce4244fb6 100644
--- a/clang/lib/CodeGen/CGCall.cpp
+++ b/clang/lib/CodeGen/CGCall.cpp
@@ -752,9 +752,8 @@ const CGFunctionInfo &CodeGenTypes::arrangeBuiltinFunctionDeclaration(
RequiredArgs::All);
}
-const CGFunctionInfo &
-CodeGenTypes::arrangeSYCLKernelCallerDeclaration(QualType resultType,
- const FunctionArgList &args) {
+const CGFunctionInfo &CodeGenTypes::arrangeDeviceKernelCallerDeclaration(
+ QualType resultType, const FunctionArgList &args) {
CanQualTypeList argTypes = getArgTypesForDeclaration(Context, args);
return arrangeLLVMFunctionInfo(GetReturnType(resultType), FnInfoOpts::None,
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index b38eb54036e60..8d67fe21367ac 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1238,7 +1238,7 @@ static llvm::Function *emitParallelOrTeamsOutlinedFunction(
CGOpenMPOutlinedRegionInfo CGInfo(*CS, ThreadIDVar, CodeGen, InnermostKind,
HasCancel, OutlinedHelperName);
CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
- return CGF.GenerateOpenMPCapturedStmtFunction(*CS, D.getBeginLoc());
+ return CGF.GenerateOpenMPCapturedStmtFunction(*CS, D);
}
std::string CGOpenMPRuntime::getOutlinedHelperName(StringRef Name) const {
@@ -6227,7 +6227,7 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
CGOpenMPTargetRegionInfo CGInfo(CS, CodeGen, EntryFnName);
CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
- return CGF.GenerateOpenMPCapturedStmtFunction(CS, D.getBeginLoc());
+ return CGF.GenerateOpenMPCapturedStmtFunction(CS, D);
};
cantFail(OMPBuilder.emitTargetRegionFunction(
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index 1360680cc6640..d72cd8fbfd608 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -471,12 +471,13 @@ struct FunctionOptions {
const StringRef FunctionName;
/// Location of the non-debug version of the outlined function.
SourceLocation Loc;
+ const bool IsDeviceKernel = false;
explicit FunctionOptions(const CapturedStmt *S, bool UIntPtrCastRequired,
bool RegisterCastedArgsOnly, StringRef FunctionName,
- SourceLocation Loc)
+ SourceLocation Loc, bool IsDeviceKernel)
: S(S), UIntPtrCastRequired(UIntPtrCastRequired),
RegisterCastedArgsOnly(UIntPtrCastRequired && RegisterCastedArgsOnly),
- FunctionName(FunctionName), Loc(Loc) {}
+ FunctionName(FunctionName), Loc(Loc), IsDeviceKernel(IsDeviceKernel) {}
};
} // namespace
@@ -570,7 +571,11 @@ static llvm::Function *emitOutlinedFunctionPrologue(
// Create the function declaration.
const CGFunctionInfo &FuncInfo =
- CGM.getTypes().arrangeBuiltinFunctionDeclaration(Ctx.VoidTy, TargetArgs);
+ FO.IsDeviceKernel
+ ? CGM.getTypes().arrangeDeviceKernelCallerDeclaration(Ctx.VoidTy,
+ TargetArgs)
+ : CGM.getTypes().arrangeBuiltinFunctionDeclaration(Ctx.VoidTy,
+ TargetArgs);
llvm::FunctionType *FuncLLVMTy = CGM.getTypes().GetFunctionType(FuncInfo);
auto *F =
@@ -664,9 +669,9 @@ static llvm::Function *emitOutlinedFunctionPrologue(
return F;
}
-llvm::Function *
-CodeGenFunction::GenerateOpenMPCapturedStmtFunction(const CapturedStmt &S,
- SourceLocation Loc) {
+llvm::Function *CodeGenFunction::GenerateOpenMPCapturedStmtFunction(
+ const CapturedStmt &S, const OMPExecutableDirective &D) {
+ SourceLocation Loc = D.getBeginLoc();
assert(
CapturedStmtInfo &&
"CapturedStmtInfo should be set when generating the captured function");
@@ -682,7 +687,10 @@ CodeGenFunction::GenerateOpenMPCapturedStmtFunction(const CapturedStmt &S,
SmallString<256> Buffer;
llvm::raw_svector_ostream Out(Buffer);
Out << CapturedStmtInfo->getHelperName();
-
+ OpenMPDirectiveKind EKind = getEffectiveDirectiveKind(D);
+ bool IsDeviceKernel = CGM.getOpenMPRuntime().isGPU() &&
+ isOpenMPTargetExecutionDirective(EKind) &&
+ D.getCapturedStmt(OMPD_target) == &S;
CodeGenFunction WrapperCGF(CGM, /*suppressNewContext=*/true);
llvm::Function *WrapperF = nullptr;
if (NeedWrapperFunction) {
@@ -690,7 +698,8 @@ CodeGenFunction::GenerateOpenMPCapturedStmtFunction(const CapturedStmt &S,
// OpenMPI-IR-Builder.
FunctionOptions WrapperFO(&S, /*UIntPtrCastRequired=*/true,
/*RegisterCastedArgsOnly=*/true,
- CapturedStmtInfo->getHelperName(), Loc);
+ CapturedStmtInfo->getHelperName(), Loc,
+ IsDeviceKernel);
WrapperCGF.CapturedStmtInfo = CapturedStmtInfo;
WrapperF =
emitOutlinedFunctionPrologue(WrapperCGF, Args, LocalAddrs, VLASizes,
@@ -698,7 +707,7 @@ CodeGenFunction::GenerateOpenMPCapturedStmtFunction(const CapturedStmt &S,
Out << "_debug__";
}
FunctionOptions FO(&S, !NeedWrapperFunction, /*RegisterCastedArgsOnly=*/false,
- Out.str(), Loc);
+ Out.str(), Loc, !NeedWrapperFunction && IsDeviceKernel);
llvm::Function *F = emitOutlinedFunctionPrologue(
*this, WrapperArgs, WrapperLocalAddrs, WrapperVLASizes, CXXThisValue, FO);
CodeGenFunction::OMPPrivateScope LocalScope(*this);
@@ -6119,13 +6128,13 @@ void CodeGenFunction::EmitOMPDistributeDirective(
emitOMPDistributeDirective(S, *this, CGM);
}
-static llvm::Function *emitOutlinedOrderedFunction(CodeGenModule &CGM,
- const CapturedStmt *S,
- SourceLocation Loc) {
+static llvm::Function *
+emitOutlinedOrderedFunction(CodeGenModule &CGM, const CapturedStmt *S,
+ const OMPExecutableDirective &D) {
CodeGenFunction CGF(CGM, /*suppressNewContext=*/true);
CodeGenFunction::CGCapturedStmtInfo CapStmtInfo;
CGF.CapturedStmtInfo = &CapStmtInfo;
- llvm::Function *Fn = CGF.GenerateOpenMPCapturedStmtFunction(*S, Loc);
+ llvm::Function *Fn = CGF.GenerateOpenMPCapturedStmtFunction(*S, D);
Fn->setDoesNotRecurse();
return Fn;
}
@@ -6190,8 +6199,7 @@ void CodeGenFunction::EmitOMPOrderedDirective(const OMPOrderedDirective &S) {
Builder, /*CreateBranch=*/false, ".ordered.after");
llvm::SmallVector<llvm::Value *, 16> CapturedVars;
GenerateOpenMPCapturedVars(*CS, CapturedVars);
- llvm::Function *OutlinedFn =
- emitOutlinedOrderedFunction(CGM, CS, S.getBeginLoc());
+ llvm::Function *OutlinedFn = emitOutlinedOrderedFunction(CGM, CS, S);
assert(S.getBeginLoc().isValid() &&
"Outlined function call location must be valid.");
ApplyDebugLocation::CreateDefaultArtificial(*this, S.getBeginLoc());
@@ -6233,8 +6241,7 @@ void CodeGenFunction::EmitOMPOrderedDirective(const OMPOrderedDirective &S) {
if (C) {
llvm::SmallVector<llvm::Value *, 16> CapturedVars;
CGF.GenerateOpenMPCapturedVars(*CS, CapturedVars);
- llvm::Function *OutlinedFn =
- emitOutlinedOrderedFunction(CGM, CS, S.getBeginLoc());
+ llvm::Function *OutlinedFn = emitOutlinedOrderedFunction(CGM, CS, S);
CGM.getOpenMPRuntime().emitOutlinedFunctionCall(CGF, S.getBeginLoc(),
OutlinedFn, CapturedVars);
} else {
diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h
index 123cb4f51f828..727487b46054f 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -3694,8 +3694,9 @@ class CodeGenFunction : public CodeGenTypeCache {
llvm::Function *EmitCapturedStmt(const CapturedStmt &S, CapturedRegionKind K);
llvm::Function *GenerateCapturedStmtFunction(const CapturedStmt &S);
Address GenerateCapturedStmtArgument(const CapturedStmt &S);
- llvm::Function *GenerateOpenMPCapturedStmtFunction(const CapturedStmt &S,
- SourceLocation Loc);
+ llvm::Function *
+ GenerateOpenMPCapturedStmtFunction(const CapturedStmt &S,
+ const OMPExecutableDirective &D);
void GenerateOpenMPCapturedVars(const CapturedStmt &S,
SmallVectorImpl<llvm::Value *> &CapturedVars);
void emitOMPSimpleStore(LValue LVal, RValue RVal, QualType RValTy,
diff --git a/clang/lib/CodeGen/CodeGenSYCL.cpp b/clang/lib/CodeGen/CodeGenSYCL.cpp
index b9a96fe8ab838..7d66d96ad0a1b 100644
--- a/clang/lib/CodeGen/CodeGenSYCL.cpp
+++ b/clang/lib/CodeGen/CodeGenSYCL.cpp
@@ -49,7 +49,7 @@ void CodeGenModule::EmitSYCLKernelCaller(const FunctionDecl *KernelEntryPointFn,
// Compute the function info and LLVM function type.
const CGFunctionInfo &FnInfo =
- getTypes().arrangeSYCLKernelCallerDeclaration(Ctx.VoidTy, Args);
+ getTypes().arrangeDeviceKernelCallerDeclaration(Ctx.VoidTy, Args);
llvm::FunctionType *FnTy = getTypes().GetFunctionType(FnInfo);
// Retrieve the generated name for the SYCL kernel caller function.
diff --git a/clang/lib/CodeGen/CodeGenTypes.h b/clang/lib/CodeGen/CodeGenTypes.h
index 29f6f1ec80bc3..9de7e0a83579d 100644
--- a/clang/lib/CodeGen/CodeGenTypes.h
+++ b/clang/lib/CodeGen/CodeGenTypes.h
@@ -229,12 +229,12 @@ class CodeGenTypes {
const CGFunctionInfo &arrangeBuiltinFunctionCall(QualType resultType,
const CallArgList &args);
- /// A SYCL kernel caller function is an offload device entry point function
+ /// A device kernel caller function is an offload device entry point function
/// with a target device dependent calling convention such as amdgpu_kernel,
/// ptx_kernel, or spir_kernel.
const CGFunctionInfo &
- arrangeSYCLKernelCallerDeclaration(QualType resultType,
- const FunctionArgList &args);
+ arrangeDeviceKernelCallerDeclaration(QualType resultType,
+ const FunctionArgList &args);
/// Objective-C methods are C functions with some implicit parameters.
const CGFunctionInfo &arrangeObjCMethodDeclaration(const ObjCMethodDecl *MD);
diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp
index 53806249ded60..01c33d1470765 100644
--- a/clang/lib/CodeGen/Targets/SPIR.cpp
+++ b/clang/lib/CodeGen/Targets/SPIR.cpp
@@ -132,10 +132,12 @@ ABIArgInfo SPIRVABIInfo::classifyReturnType(QualType RetTy) const {
}
ABIArgInfo SPIRVABIInfo::classifyKernelArgumentType(QualType Ty) const {
- if (getContext().getLangOpts().CUDAIsDevice) {
+ if (getContext().getLangOpts().CUDAIsDevice ||
+ getContext().getLangOpts().OpenMPIsTargetDevice) {
// Coerce pointer arguments with default address space to CrossWorkGroup
- // pointers for HIPSPV/CUDASPV. When the language mode is HIP/CUDA, the
- // SPIRTargetInfo maps cuda_device to SPIR-V's CrossWorkGroup address space.
+ // pointers for HIPSPV/CUDASPV/OMPSPV. When the language mode is
+ // HIP/CUDA/OMP, the SPIRTargetInfo maps cuda_device to SPIR-V's
+ // CrossWorkGroup address space.
llvm::Type *LTy = CGT.ConvertType(Ty);
auto DefaultAS = getContext().getTargetAddressSpace(LangAS::Default);
auto GlobalAS = getContext().getTargetAddressSpace(LangAS::cuda_device);
diff --git a/clang/test/OpenMP/spirv_kernel_addrspace.cpp b/clang/test/OpenMP/spirv_kernel_addrspace.cpp
new file mode 100644
index 0000000000000..cea7e9958c341
--- /dev/null
+++ b/clang/test/OpenMP/spirv_kernel_addrspace.cpp
@@ -0,0 +1,24 @@
+// 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
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=spirv64-intel -emit-llvm-bc %s -o %t-host.bc -DTEAMS
+// 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 -DTEAMS -o - | FileCheck %s
+// expected-no-diagnostics
+
+// CHECK: define weak_odr protected spir_kernel void @__omp_offloading_{{.*}}(ptr addrspace(1) noalias noundef %{{.*}}, ptr addrspace(1) noundef align 4 dereferenceable(128) %{{.*}})
+
+int main() {
+ int x[32] = {0};
+
+#ifdef TEAMS
+#pragma omp target teams
+#else
+#pragma omp target
+#endif
+ for(int i = 0; i < 32; i++) {
+ if(i > 0)
+ x[i] = x[i-1] + i;
+ }
+
+return x[31];
+}
+
|
jhuber6
reviewed
Sep 8, 2025
clang/lib/CodeGen/CGStmtOpenMP.cpp
Outdated
Contributor
There was a problem hiding this comment.
It's tough to tell outside of the LLVM-IR because we generate many things to kernels. What's here seems to make sense to me, but I'm assuming that the check on target applies to all the variants of target.
clang/lib/CodeGen/Targets/SPIR.cpp
Outdated
Comment on lines
135
to
136
| if (getContext().getLangOpts().CUDAIsDevice || | ||
| getContext().getLangOpts().OpenMPIsTargetDevice) { |
Contributor
There was a problem hiding this comment.
Don't we have a helper for all the different languages in device mode? I could be wrong since I can't remember what it'd be called.
Member
Author
There was a problem hiding this comment.
Yeah and it's even more sad because I'm the one who added that function :P
Fixed in latest commit and reworded comment to remove irrelevant part and be more clear.
Member
Author
There was a problem hiding this comment.
Sorry that part of the comment I removed was not only relevant but exposed a possible issue, fixed in latest commit.
Signed-off-by: Sarnie, Nick <[email protected]>
jhuber6
approved these changes
Sep 9, 2025
Signed-off-by: Sarnie, Nick <[email protected]>
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
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.
In SPIR-V, kernel arguments are not allowed to be in the Generic AS, in both Intel's internal SPIR-V offloading implementation as well as HIPSPV,
CrossWorkgroupAS1 is used. Do the same for OMPSPV.Currently with Generic AS the
llvm-spirvtranslator blows up if we are using it, and if not, the GPU runtime blows up.To get the existing logic to set the correct AS to kick in, we need to know if the function is a kernel or not at the time we first create the function that may end up as the kernel.
I use the existing
arrangeSYCLKernelCallerDeclarationfunction to do the right kernel ABI computation, but since the function is not specific to SYCL anymore because I merged all the device kernel clang attributes into one.Rename the function to be accurate to the current behavior,
arrangeDeviceKernelCallerDeclaration.