[clang] [CIR][SPIR-V] Set spir_kernel calling convention for AMDGCN-flavored HIP kernels (PR #214246)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Wed Aug 26 03:02:14 PDT 2026
https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/214246
>From e0289b4451833784512fa7f6f3702ea10802e536 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Wed, 5 Aug 2026 17:05:13 +0200
Subject: [PATCH 1/5] [CIR][SPIR-V] Set spir_kernel calling convention for
AMDGCN-flavored SPIR-V HIP kernels
---
clang/lib/CIR/CodeGen/Targets/SPIRV.cpp | 9 +++++++-
.../CIR/CodeGenHIP/amdgcnspirv-kernel.hip | 23 +++++++++++++++++++
2 files changed, 31 insertions(+), 1 deletion(-)
create mode 100644 clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip
diff --git a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
index 643c635128d09..b56bbfe41e87c 100644
--- a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
+++ b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
@@ -42,11 +42,18 @@ class SPIRVTargetCIRGenInfo : public TargetCIRGenInfo {
if (!fd)
return;
+ auto func = mlir::cast<cir::FuncOp>(global);
+
if (cgm.getLangOpts().OpenCL &&
DeviceKernelAttr::isOpenCLSpelling(fd->getAttr<DeviceKernelAttr>())) {
- auto func = mlir::cast<cir::FuncOp>(global);
func.setCallingConv(cir::CallingConv::SpirKernel);
+ return;
}
+
+ if (cgm.getLangOpts().HIP &&
+ cgm.getTriple().getVendor() == llvm::Triple::AMD &&
+ fd->hasAttr<CUDAGlobalAttr>())
+ func.setCallingConv(cir::CallingConv::SpirKernel);
}
};
diff --git a/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip b/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip
new file mode 100644
index 0000000000000..039ab35f1c906
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip
@@ -0,0 +1,23 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fclangir \
+// RUN: -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.cir
+
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fclangir \
+// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.ll
+
+// Test that HIP kernels on AMDGCN-flavored SPIR-V get the spir_kernel
+// calling convention.
+
+#define __global__ __attribute__((global))
+#define __device__ __attribute__((device))
+
+// CIR: cir.func{{.*}} @_Z13kernel_scalari{{.*}} cc(spir_kernel)
+// LLVM: define spir_kernel void @_Z13kernel_scalari
+__global__ void kernel_scalar(int a) {}
+
+// CIR: cir.func{{.*}} @_Z9device_fni
+// CIR-NOT: cc(spir_kernel)
+// LLVM: define{{.*}} void @_Z9device_fni
+__device__ void device_fn(int a) {}
>From 5a17e67fd7786430df8330a9b1c36535984a3ec2 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Wed, 5 Aug 2026 17:42:59 +0200
Subject: [PATCH 2/5] fix integerattr
The program address space DLTI entry is stored as an unsigned integer
attribute, but was read with getInt(), which asserts on signless types.
---
clang/lib/CIR/Dialect/IR/CIRDataLayout.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/lib/CIR/Dialect/IR/CIRDataLayout.cpp b/clang/lib/CIR/Dialect/IR/CIRDataLayout.cpp
index d7d563d79a0fa..5ce0c8f8e7a2c 100644
--- a/clang/lib/CIR/Dialect/IR/CIRDataLayout.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRDataLayout.cpp
@@ -26,7 +26,7 @@ void CIRDataLayout::reset(mlir::DataLayoutSpecInterface spec) {
if (mlir::DataLayoutEntryInterface entry =
spec.getSpecForIdentifier(addrSpKey))
if (auto val = llvm::dyn_cast<mlir::IntegerAttr>(entry.getValue()))
- programAddrSpace = val.getInt();
+ programAddrSpace = val.getUInt();
}
}
>From 93eb6b0d66627648bc57bbee5e9d49c58e5cc2da Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Mon, 24 Aug 2026 12:43:49 +0200
Subject: [PATCH 3/5] rm AMD triple check
---
clang/lib/CIR/CodeGen/Targets/SPIRV.cpp | 4 +---
1 file changed, 1 insertion(+), 3 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
index b56bbfe41e87c..8a19fa1f8a9c0 100644
--- a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
+++ b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
@@ -50,9 +50,7 @@ class SPIRVTargetCIRGenInfo : public TargetCIRGenInfo {
return;
}
- if (cgm.getLangOpts().HIP &&
- cgm.getTriple().getVendor() == llvm::Triple::AMD &&
- fd->hasAttr<CUDAGlobalAttr>())
+ if (cgm.getLangOpts().HIP && fd->hasAttr<CUDAGlobalAttr>())
func.setCallingConv(cir::CallingConv::SpirKernel);
}
};
>From d22d3b3b8b0487ac731153ddda2dc2df2485c4a4 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Wed, 26 Aug 2026 08:36:39 +0200
Subject: [PATCH 4/5] Align CI and classic codegen
---
clang/lib/CIR/CodeGen/CIRGenCall.cpp | 32 ++++++++++++++++------
clang/lib/CIR/CodeGen/CIRGenFunctionInfo.h | 21 ++++++++++++--
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 6 +---
clang/lib/CIR/CodeGen/CIRGenTypes.cpp | 19 +++++++++++--
clang/lib/CIR/CodeGen/CIRGenTypes.h | 3 ++
clang/lib/CIR/CodeGen/TargetInfo.cpp | 8 ++++++
clang/lib/CIR/CodeGen/TargetInfo.h | 6 ++++
clang/lib/CIR/CodeGen/Targets/SPIRV.cpp | 29 ++++++--------------
8 files changed, 85 insertions(+), 39 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenCall.cpp b/clang/lib/CIR/CodeGen/CIRGenCall.cpp
index 28670cf31694a..bf94d0fcbeb24 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCall.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenCall.cpp
@@ -26,9 +26,11 @@
using namespace clang;
using namespace clang::CIRGen;
-CIRGenFunctionInfo *CIRGenFunctionInfo::create(
- FunctionType::ExtInfo info, bool isInstanceMethod, CanQualType resultType,
- llvm::ArrayRef<CanQualType> argTypes, RequiredArgs required) {
+CIRGenFunctionInfo *
+CIRGenFunctionInfo::create(cir::CallingConv cirCC, FunctionType::ExtInfo info,
+ bool isInstanceMethod, CanQualType resultType,
+ llvm::ArrayRef<CanQualType> argTypes,
+ RequiredArgs required) {
// The first slot allocated for arg type slot is for the return value.
void *buffer = operator new(
totalSizeToAlloc<CanQualType>(argTypes.size() + 1));
@@ -37,6 +39,8 @@ CIRGenFunctionInfo *CIRGenFunctionInfo::create(
CIRGenFunctionInfo *fi = new (buffer) CIRGenFunctionInfo();
+ fi->callingConvention = llvm::to_underlying(cirCC);
+ fi->astCallingConvention = info.getCC();
fi->noReturn = info.getNoReturn();
fi->instanceMethod = isInstanceMethod;
@@ -311,7 +315,7 @@ void CIRGenModule::constructAttributeList(
llvm::MutableArrayRef<mlir::NamedAttrList> argAttrs,
mlir::NamedAttrList &retAttrs, cir::CallingConv &callingConv,
cir::SideEffect &sideEffect, bool attrOnCallSite, bool isThunk) {
- assert(!cir::MissingFeatures::opCallCallConv());
+ callingConv = info.getCallingConvention();
sideEffect = cir::SideEffect::All;
auto addUnitAttr = [&](llvm::StringRef name) {
@@ -1028,6 +1032,17 @@ CIRGenTypes::arrangeBuiltinFunctionCall(QualType resultType,
FunctionType::ExtInfo(), RequiredArgs::All);
}
+/// Set calling convention for CUDA/HIP kernel.
+static void setCUDAKernelCallingConvention(CanQualType &funcTy,
+ CIRGenModule &cgm,
+ const FunctionDecl *fd) {
+ if (fd->hasAttr<CUDAGlobalAttr>()) {
+ const FunctionType *ft = funcTy->getAs<FunctionType>();
+ cgm.getTargetCIRGenInfo().setCUDAKernelCallingConvention(ft);
+ funcTy = ft->getCanonicalTypeUnqualified();
+ }
+}
+
/// Arrange the argument and result information for a declaration or definition
/// of the given C++ non-static member function. The member function must be an
/// ordinary function, i.e. not a constructor or destructor.
@@ -1036,9 +1051,9 @@ CIRGenTypes::arrangeCXXMethodDeclaration(const CXXMethodDecl *md) {
assert(!isa<CXXConstructorDecl>(md) && "wrong method for constructors!");
assert(!isa<CXXDestructorDecl>(md) && "wrong method for destructors!");
- auto prototype =
- md->getType()->getCanonicalTypeUnqualified().getAs<FunctionProtoType>();
- assert(!cir::MissingFeatures::cudaSupport());
+ CanQualType funcTy = md->getType()->getCanonicalTypeUnqualified();
+ setCUDAKernelCallingConvention(funcTy, cgm, md);
+ auto prototype = funcTy.getAs<FunctionProtoType>();
// Mirrors classic CodeGen's check at CGCall.cpp. C++23 explicit-object
// member functions (P0847R7, `void f(this Self&&)`) do not receive an
@@ -1088,8 +1103,7 @@ CIRGenTypes::arrangeFunctionDeclaration(const FunctionDecl *fd) {
CanQualType funcTy = fd->getType()->getCanonicalTypeUnqualified();
assert(isa<FunctionType>(funcTy));
- // TODO: setCUDAKernelCallingConvention
- assert(!cir::MissingFeatures::cudaSupport());
+ setCUDAKernelCallingConvention(funcTy, cgm, fd);
// When declaring a function without a prototype, always use a non-variadic
// type.
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunctionInfo.h b/clang/lib/CIR/CodeGen/CIRGenFunctionInfo.h
index d37a6149bcafa..2957293c83623 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunctionInfo.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunctionInfo.h
@@ -17,6 +17,7 @@
#include "clang/AST/CanonicalType.h"
#include "clang/CIR/ABIArgInfo.h"
+#include "clang/CIR/Dialect/IR/CIROpsEnums.h"
#include "clang/CIR/MissingFeatures.h"
#include "llvm/ADT/FoldingSet.h"
#include "llvm/Support/TrailingObjects.h"
@@ -84,6 +85,12 @@ class RequiredArgs {
class CIRGenFunctionInfo final
: public llvm::FoldingSetNode,
private llvm::TrailingObjects<CIRGenFunctionInfo, CanQualType> {
+ /// The CIR-level calling convention to use for this function.
+ unsigned callingConvention : 8;
+
+ /// The AST-level calling convention this function was declared with.
+ unsigned astCallingConvention : 8;
+
// Whether this function has noreturn.
LLVM_PREFERRED_TYPE(bool)
unsigned noReturn : 1;
@@ -107,13 +114,14 @@ class CIRGenFunctionInfo final
// here instead of explicit false/0.
return FunctionType::ExtInfo(
isNoReturn(), /*getHasRegParm=*/false, /*getRegParm=*/false,
- /*getASTCallingConvention=*/CallingConv(0), /*isReturnsRetained=*/false,
+ getASTCallingConvention(), /*isReturnsRetained=*/false,
/*isNoCallerSavedRegs=*/false, /*isNoCfCheck=*/false,
/*isCmseNSCall=*/false);
}
public:
- static CIRGenFunctionInfo *create(FunctionType::ExtInfo info,
+ static CIRGenFunctionInfo *create(cir::CallingConv cirCC,
+ FunctionType::ExtInfo info,
bool instanceMethod, CanQualType resultType,
llvm::ArrayRef<CanQualType> argTypes,
RequiredArgs required);
@@ -133,6 +141,7 @@ class CIRGenFunctionInfo final
FunctionType::ExtInfo info, RequiredArgs required,
CanQualType resultType,
llvm::ArrayRef<CanQualType> argTypes) {
+ id.AddInteger(info.getCC());
id.AddBoolean(instanceMethod);
id.AddBoolean(info.getNoReturn());
id.AddInteger(required.getOpaqueData());
@@ -192,6 +201,14 @@ class CIRGenFunctionInfo final
bool isNoReturn() const { return noReturn; }
bool isInstanceMethod() const { return instanceMethod; }
+
+ cir::CallingConv getCallingConvention() const {
+ return static_cast<cir::CallingConv>(callingConvention);
+ }
+
+ CallingConv getASTCallingConvention() const {
+ return static_cast<CallingConv>(astCallingConvention);
+ }
};
} // namespace clang::CIRGen
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index a7143408c9bee..c1b5f1faf16d9 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -3125,11 +3125,7 @@ void CIRGenModule::setCIRFunctionAttributes(GlobalDecl globalDecl,
// TODO(cir): Check X86_VectorCall incompatibility wiht WinARM64EC
- // TODO(cir): Set the calling convention computed by constructAttributeList
- // on the function. FuncOp supports calling_conv, but target-specific
- // CodeGen is needed to set it correctly (e.g., AMDGPU kernel functions
- // should be marked with AMDGPUKernel).
- assert(!cir::MissingFeatures::opFuncCallingConv());
+ func.setCallingConv(callingConv);
}
void CIRGenModule::setFunctionAttributes(GlobalDecl globalDecl,
diff --git a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp b/clang/lib/CIR/CodeGen/CIRGenTypes.cpp
index e5af4eec7720f..469f052c66780 100644
--- a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenTypes.cpp
@@ -750,6 +750,19 @@ bool CIRGenTypes::isZeroInitializable(const RecordDecl *rd) {
return getCIRGenRecordLayout(rd).isZeroInitializable();
}
+cir::CallingConv
+CIRGenTypes::clangCallConvToCIRCallConv(clang::CallingConv cc) {
+ switch (cc) {
+ default:
+ // TODO(cir): Support the remaining target-specific calling conventions.
+ return cir::CallingConv::C;
+ case CC_SpirFunction:
+ return cir::CallingConv::SpirFunction;
+ case CC_DeviceKernel:
+ return cgm.getTargetCIRGenInfo().getDeviceKernelCallingConv();
+ }
+}
+
const CIRGenFunctionInfo &CIRGenTypes::arrangeCIRFunctionInfo(
CanQualType returnType, bool isInstanceMethod,
llvm::ArrayRef<CanQualType> argTypes, FunctionType::ExtInfo info,
@@ -773,11 +786,11 @@ const CIRGenFunctionInfo &CIRGenTypes::arrangeCIRFunctionInfo(
return *fi;
}
- assert(!cir::MissingFeatures::opCallCallConv());
+ cir::CallingConv cirCC = clangCallConvToCIRCallConv(info.getCC());
// Construction the function info. We co-allocate the ArgInfos.
- fi = CIRGenFunctionInfo::create(info, isInstanceMethod, returnType, argTypes,
- required);
+ fi = CIRGenFunctionInfo::create(cirCC, info, isInstanceMethod, returnType,
+ argTypes, required);
functionInfos.InsertNode(fi, insertPos);
return *fi;
diff --git a/clang/lib/CIR/CodeGen/CIRGenTypes.h b/clang/lib/CIR/CodeGen/CIRGenTypes.h
index a7827f76bd5f2..4c9280876882d 100644
--- a/clang/lib/CIR/CodeGen/CIRGenTypes.h
+++ b/clang/lib/CIR/CodeGen/CIRGenTypes.h
@@ -219,6 +219,9 @@ class CIRGenTypes {
const CIRGenFunctionInfo &arrangeFreeFunctionCall(const CallArgList &args,
const FunctionType *fnType);
+ /// Convert a clang calling convention to a CIR calling convention.
+ cir::CallingConv clangCallConvToCIRCallConv(clang::CallingConv cc);
+
const CIRGenFunctionInfo &
arrangeCIRFunctionInfo(CanQualType returnType, bool isInstanceMethod,
llvm::ArrayRef<CanQualType> argTypes,
diff --git a/clang/lib/CIR/CodeGen/TargetInfo.cpp b/clang/lib/CIR/CodeGen/TargetInfo.cpp
index ba7eeb29dd252..42906a5f68c12 100644
--- a/clang/lib/CIR/CodeGen/TargetInfo.cpp
+++ b/clang/lib/CIR/CodeGen/TargetInfo.cpp
@@ -143,6 +143,14 @@ bool TargetCIRGenInfo::isNoProtoCallVariadic(
return false;
}
+cir::CallingConv TargetCIRGenInfo::getDeviceKernelCallingConv() const {
+ // Device kernels are entered through a runtime API, not called as normal
+ // sub-functions, so a modified C calling convention is used.
+ assert(getABIInfo().cgt.getASTContext().getLangOpts().OpenCL &&
+ "Kernel calling convention only defined for OpenCL");
+ return cir::CallingConv::C;
+}
+
clang::LangAS
TargetCIRGenInfo::getGlobalVarAddressSpace(CIRGenModule &cgm,
const clang::VarDecl *d) const {
diff --git a/clang/lib/CIR/CodeGen/TargetInfo.h b/clang/lib/CIR/CodeGen/TargetInfo.h
index 308d472234f99..326a05bf22e86 100644
--- a/clang/lib/CIR/CodeGen/TargetInfo.h
+++ b/clang/lib/CIR/CodeGen/TargetInfo.h
@@ -128,6 +128,12 @@ class TargetCIRGenInfo {
return false;
}
+ /// Returns the calling convention used for device kernels on this target.
+ virtual cir::CallingConv getDeviceKernelCallingConv() const;
+
+ virtual void
+ setCUDAKernelCallingConvention(const clang::FunctionType *&ft) const {}
+
/// Corrects the MLIR type for a given constraint and "usual"
/// type.
///
diff --git a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
index 8a19fa1f8a9c0..598cd693d6d26 100644
--- a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
+++ b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
@@ -32,26 +32,15 @@ class SPIRVTargetCIRGenInfo : public TargetCIRGenInfo {
SPIRVTargetCIRGenInfo(CIRGenTypes &cgt)
: TargetCIRGenInfo(std::make_unique<SPIRVABIInfo>(cgt)) {}
- void setTargetAttributes(const clang::Decl *decl, mlir::Operation *global,
- CIRGenModule &cgm) const override {
- auto globalValue = mlir::cast<cir::CIRGlobalValueInterface>(global);
- if (globalValue.isDeclaration())
- return;
-
- const auto *fd = dyn_cast_or_null<FunctionDecl>(decl);
- if (!fd)
- return;
-
- auto func = mlir::cast<cir::FuncOp>(global);
-
- if (cgm.getLangOpts().OpenCL &&
- DeviceKernelAttr::isOpenCLSpelling(fd->getAttr<DeviceKernelAttr>())) {
- func.setCallingConv(cir::CallingConv::SpirKernel);
- return;
- }
-
- if (cgm.getLangOpts().HIP && fd->hasAttr<CUDAGlobalAttr>())
- func.setCallingConv(cir::CallingConv::SpirKernel);
+ cir::CallingConv getDeviceKernelCallingConv() const override {
+ return cir::CallingConv::SpirKernel;
+ }
+
+ void setCUDAKernelCallingConvention(const FunctionType *&ft) const override {
+ // Convert HIP kernels to SPIR-V kernels.
+ if (getABIInfo().cgt.getASTContext().getLangOpts().HIP)
+ ft = getABIInfo().cgt.getASTContext().adjustFunctionType(
+ ft, ft->getExtInfo().withCallingConv(CC_DeviceKernel));
}
};
>From b8313e894cde477279e7100b14e6654979f6b882 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Wed, 26 Aug 2026 12:01:46 +0200
Subject: [PATCH 5/5] fix build
---
clang/lib/CIR/CodeGen/CIRGenTypes.cpp | 11 +++++++----
clang/lib/CIR/CodeGen/TargetInfo.h | 6 ------
clang/lib/CIR/CodeGen/Targets/SPIRV.cpp | 4 ----
3 files changed, 7 insertions(+), 14 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp b/clang/lib/CIR/CodeGen/CIRGenTypes.cpp
index a4982ff707cf1..f9fa018d70540 100644
--- a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenTypes.cpp
@@ -756,13 +756,16 @@ bool CIRGenTypes::isZeroInitializable(const RecordDecl *rd) {
cir::CallingConv
CIRGenTypes::clangCallConvToCIRCallConv(clang::CallingConv cc) {
switch (cc) {
- default:
- // TODO(cir): Support the remaining target-specific calling conventions.
+ case CC_C:
+ // SPIR/SPIR-V lowers the default CC to spir_func, not plain C.
+ if (cgm.getTriple().isSPIROrSPIRV())
+ return cir::CallingConv::SpirFunction;
return cir::CallingConv::C;
- case CC_SpirFunction:
- return cir::CallingConv::SpirFunction;
case CC_DeviceKernel:
return cgm.getTargetCIRGenInfo().getDeviceKernelCallingConv();
+ default:
+ // TODO(cir): Support the remaining target-specific calling conventions.
+ return cir::CallingConv::C;
}
}
diff --git a/clang/lib/CIR/CodeGen/TargetInfo.h b/clang/lib/CIR/CodeGen/TargetInfo.h
index 54cb256bcb774..0e91f4e95a4a9 100644
--- a/clang/lib/CIR/CodeGen/TargetInfo.h
+++ b/clang/lib/CIR/CodeGen/TargetInfo.h
@@ -153,12 +153,6 @@ class TargetCIRGenInfo {
mlir::Operation *global,
CIRGenModule &module) const {}
- /// Get the CIR calling convention to use for a device kernel entry point
- /// (e.g. an OpenCL/SYCL or CUDA/HIP kernel) on this target.
- virtual cir::CallingConv getDeviceKernelCallingConv() const {
- return cir::CallingConv::C;
- }
-
virtual bool isScalarizableAsmOperand(CIRGenFunction &cgf,
mlir::Type ty) const {
return false;
diff --git a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
index e88086242e99e..598cd693d6d26 100644
--- a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
+++ b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
@@ -42,10 +42,6 @@ class SPIRVTargetCIRGenInfo : public TargetCIRGenInfo {
ft = getABIInfo().cgt.getASTContext().adjustFunctionType(
ft, ft->getExtInfo().withCallingConv(CC_DeviceKernel));
}
-
- cir::CallingConv getDeviceKernelCallingConv() const override {
- return cir::CallingConv::SpirKernel;
- }
};
} // namespace
More information about the cfe-commits
mailing list