[clang] e96c3fc - [CIR][SPIR-V] Set spir_kernel calling convention for AMDGCN-flavored HIP kernels (#214246)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Sep 4 04:58:47 PDT 2026
Author: Arseniy Obolenskiy
Date: 2026-09-04T13:58:42+02:00
New Revision: e96c3fcd9824a9ae8dfe6207188a0c9fa4f65e90
URL: https://github.com/llvm/llvm-project/commit/e96c3fcd9824a9ae8dfe6207188a0c9fa4f65e90
DIFF: https://github.com/llvm/llvm-project/commit/e96c3fcd9824a9ae8dfe6207188a0c9fa4f65e90.diff
LOG: [CIR][SPIR-V] Set spir_kernel calling convention for AMDGCN-flavored HIP kernels (#214246)
CIR emitted no calling convention for HIP `__global__` kernels on the
`spirv64-amd-amdhsa` target, unlike generic SPIR-V, which already gets
`spirv_kernel` CC
Added:
clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip
Modified:
clang/lib/CIR/CodeGen/CIRGenCall.cpp
clang/lib/CIR/CodeGen/CIRGenFunctionInfo.h
clang/lib/CIR/CodeGen/CIRGenModule.cpp
clang/lib/CIR/CodeGen/CIRGenTypes.cpp
clang/lib/CIR/CodeGen/CIRGenTypes.h
clang/lib/CIR/CodeGen/TargetInfo.cpp
clang/lib/CIR/CodeGen/TargetInfo.h
clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
clang/lib/CIR/Dialect/IR/CIRDataLayout.cpp
clang/test/CIR/CodeGenOpenCL/address-space-local-var.clcpp
Removed:
################################################################################
diff --git a/clang/lib/CIR/CodeGen/CIRGenCall.cpp b/clang/lib/CIR/CodeGen/CIRGenCall.cpp
index 3e689e031f7ad..3a4b7cecf2e08 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;
@@ -319,7 +323,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) {
@@ -1047,6 +1051,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.
@@ -1055,9 +1070,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
@@ -1107,8 +1122,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 91f22b2fc5ca1..b70f4c0c0b043 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -332,6 +332,8 @@ const TargetCIRGenInfo &CIRGenModule::getTargetCIRGenInfo() {
theTargetCIRGenInfo = createAMDGPUTargetCIRGenInfo(genTypes);
return *theTargetCIRGenInfo;
}
+ case llvm::Triple::spir:
+ case llvm::Triple::spir64:
case llvm::Triple::spirv:
case llvm::Triple::spirv32:
case llvm::Triple::spirv64:
@@ -3231,11 +3233,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 2e80a84a4bd20..f2cb875e908a7 100644
--- a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenTypes.cpp
@@ -757,6 +757,22 @@ bool CIRGenTypes::isZeroInitializable(const RecordDecl *rd) {
return getCIRGenRecordLayout(rd).isZeroInitializable();
}
+cir::CallingConv
+CIRGenTypes::clangCallConvToCIRCallConv(clang::CallingConv cc) {
+ switch (cc) {
+ 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_DeviceKernel:
+ return cgm.getTargetCIRGenInfo().getDeviceKernelCallingConv();
+ default:
+ // TODO(cir): Support the remaining target-specific calling conventions.
+ return cir::CallingConv::C;
+ }
+}
+
const CIRGenFunctionInfo &CIRGenTypes::arrangeCIRFunctionInfo(
CanQualType returnType, bool isInstanceMethod,
llvm::ArrayRef<CanQualType> argTypes, FunctionType::ExtInfo info,
@@ -780,11 +796,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.insert(fi, insertToken);
return *fi;
diff --git a/clang/lib/CIR/CodeGen/CIRGenTypes.h b/clang/lib/CIR/CodeGen/CIRGenTypes.h
index c5f8b521ed888..e2045f9e58d04 100644
--- a/clang/lib/CIR/CodeGen/CIRGenTypes.h
+++ b/clang/lib/CIR/CodeGen/CIRGenTypes.h
@@ -226,6 +226,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 51c9cfc056eee..8dd6147c37f0b 100644
--- a/clang/lib/CIR/CodeGen/TargetInfo.cpp
+++ b/clang/lib/CIR/CodeGen/TargetInfo.cpp
@@ -199,6 +199,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 4f636ab597201..7f6363b6a5176 100644
--- a/clang/lib/CIR/CodeGen/TargetInfo.h
+++ b/clang/lib/CIR/CodeGen/TargetInfo.h
@@ -157,17 +157,17 @@ 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;
}
+ /// 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 f2d9810b36061..598cd693d6d26 100644
--- a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
+++ b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
@@ -32,26 +32,16 @@ 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;
-
- if (cgm.getLangOpts().OpenCL &&
- DeviceKernelAttr::isOpenCLSpelling(fd->getAttr<DeviceKernelAttr>())) {
- auto func = mlir::cast<cir::FuncOp>(global);
- 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));
+ }
};
} // namespace
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();
}
}
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) {}
diff --git a/clang/test/CIR/CodeGenOpenCL/address-space-local-var.clcpp b/clang/test/CIR/CodeGenOpenCL/address-space-local-var.clcpp
index bfce737207e87..84dea3b6daead 100644
--- a/clang/test/CIR/CodeGenOpenCL/address-space-local-var.clcpp
+++ b/clang/test/CIR/CodeGenOpenCL/address-space-local-var.clcpp
@@ -21,7 +21,7 @@
// CIR: %[[CAST:.*]] = cir.cast address_space %[[GR_VAL]] : !cir.ptr<!s32i> -> !cir.ptr<!s32i>
// CIR: cir.store {{.*}} %[[CAST]], %[[R]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
-// LLVM: define dso_local void @k(ptr noundef %[[ARG:.*]])
+// LLVM: define dso_local spir_kernel void @k(ptr noundef %[[ARG:.*]])
// LLVM: %[[GP_ADDR:.*]] = alloca ptr
// LLVM: %[[GR_ADDR:.*]] = alloca ptr
// LLVM: %[[R_ADDR:.*]] = alloca ptr
More information about the cfe-commits
mailing list