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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
