r328795 - Set calling convention for CUDA kernel
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Thu Mar 29 08:02:08 PDT 2018
Author: yaxunl
Date: Thu Mar 29 08:02:08 2018
New Revision: 328795
URL: http://llvm.org/viewvc/llvm-project?rev=328795&view=rev
Log:
Set calling convention for CUDA kernel
This patch sets target specific calling convention for CUDA kernels in IR.
Patch by Greg Rodgers.
Revised and lit test added by Yaxun Liu.
Differential Revision: https://reviews.llvm.org/D44747
Added:
cfe/trunk/test/CodeGenCUDA/kernel-amdgcn.cu
Modified:
cfe/trunk/include/clang/Basic/Specifiers.h
cfe/trunk/lib/AST/ItaniumMangle.cpp
cfe/trunk/lib/AST/Type.cpp
cfe/trunk/lib/AST/TypePrinter.cpp
cfe/trunk/lib/CodeGen/CGCall.cpp
cfe/trunk/lib/CodeGen/CGDebugInfo.cpp
cfe/trunk/lib/CodeGen/TargetInfo.cpp
cfe/trunk/lib/CodeGen/TargetInfo.h
cfe/trunk/lib/Sema/SemaExpr.cpp
cfe/trunk/lib/Sema/SemaOverload.cpp
cfe/trunk/lib/Sema/SemaType.cpp
cfe/trunk/tools/libclang/CXType.cpp
Modified: cfe/trunk/include/clang/Basic/Specifiers.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/Specifiers.h?rev=328795&r1=328794&r2=328795&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/Specifiers.h (original)
+++ cfe/trunk/include/clang/Basic/Specifiers.h Thu Mar 29 08:02:08 2018
@@ -231,23 +231,24 @@ namespace clang {
/// \brief CallingConv - Specifies the calling convention that a function uses.
enum CallingConv {
- CC_C, // __attribute__((cdecl))
- CC_X86StdCall, // __attribute__((stdcall))
- CC_X86FastCall, // __attribute__((fastcall))
- CC_X86ThisCall, // __attribute__((thiscall))
+ CC_C, // __attribute__((cdecl))
+ CC_X86StdCall, // __attribute__((stdcall))
+ CC_X86FastCall, // __attribute__((fastcall))
+ CC_X86ThisCall, // __attribute__((thiscall))
CC_X86VectorCall, // __attribute__((vectorcall))
- CC_X86Pascal, // __attribute__((pascal))
- CC_Win64, // __attribute__((ms_abi))
- CC_X86_64SysV, // __attribute__((sysv_abi))
- CC_X86RegCall, // __attribute__((regcall))
- CC_AAPCS, // __attribute__((pcs("aapcs")))
- CC_AAPCS_VFP, // __attribute__((pcs("aapcs-vfp")))
- CC_IntelOclBicc, // __attribute__((intel_ocl_bicc))
- CC_SpirFunction, // default for OpenCL functions on SPIR target
- CC_OpenCLKernel, // inferred for OpenCL kernels
- CC_Swift, // __attribute__((swiftcall))
- CC_PreserveMost, // __attribute__((preserve_most))
- CC_PreserveAll, // __attribute__((preserve_all))
+ CC_X86Pascal, // __attribute__((pascal))
+ CC_Win64, // __attribute__((ms_abi))
+ CC_X86_64SysV, // __attribute__((sysv_abi))
+ CC_X86RegCall, // __attribute__((regcall))
+ CC_AAPCS, // __attribute__((pcs("aapcs")))
+ CC_AAPCS_VFP, // __attribute__((pcs("aapcs-vfp")))
+ CC_IntelOclBicc, // __attribute__((intel_ocl_bicc))
+ CC_SpirFunction, // default for OpenCL functions on SPIR target
+ CC_OpenCLKernel, // inferred for OpenCL kernels
+ CC_Swift, // __attribute__((swiftcall))
+ CC_PreserveMost, // __attribute__((preserve_most))
+ CC_PreserveAll, // __attribute__((preserve_all))
+ CC_CUDAKernel, // inferred for CUDA kernels
};
/// \brief Checks whether the given calling convention supports variadic
Modified: cfe/trunk/lib/AST/ItaniumMangle.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ItaniumMangle.cpp?rev=328795&r1=328794&r2=328795&view=diff
==============================================================================
--- cfe/trunk/lib/AST/ItaniumMangle.cpp (original)
+++ cfe/trunk/lib/AST/ItaniumMangle.cpp Thu Mar 29 08:02:08 2018
@@ -2628,6 +2628,7 @@ StringRef CXXNameMangler::getCallingConv
case CC_OpenCLKernel:
case CC_PreserveMost:
case CC_PreserveAll:
+ case CC_CUDAKernel:
// FIXME: we should be mangling all of the above.
return "";
Modified: cfe/trunk/lib/AST/Type.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/Type.cpp?rev=328795&r1=328794&r2=328795&view=diff
==============================================================================
--- cfe/trunk/lib/AST/Type.cpp (original)
+++ cfe/trunk/lib/AST/Type.cpp Thu Mar 29 08:02:08 2018
@@ -2752,6 +2752,7 @@ StringRef FunctionType::getNameForCallCo
case CC_Swift: return "swiftcall";
case CC_PreserveMost: return "preserve_most";
case CC_PreserveAll: return "preserve_all";
+ case CC_CUDAKernel: return "cuda_kernel";
}
llvm_unreachable("Invalid calling convention.");
Modified: cfe/trunk/lib/AST/TypePrinter.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/TypePrinter.cpp?rev=328795&r1=328794&r2=328795&view=diff
==============================================================================
--- cfe/trunk/lib/AST/TypePrinter.cpp (original)
+++ cfe/trunk/lib/AST/TypePrinter.cpp Thu Mar 29 08:02:08 2018
@@ -780,6 +780,10 @@ void TypePrinter::printFunctionAfter(con
case CC_OpenCLKernel:
// Do nothing. These CCs are not available as attributes.
break;
+ case CC_CUDAKernel:
+ // ToDo: print this before the function.
+ OS << " __global__";
+ break;
case CC_Swift:
OS << " __attribute__((swiftcall))";
break;
Modified: cfe/trunk/lib/CodeGen/CGCall.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGCall.cpp?rev=328795&r1=328794&r2=328795&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGCall.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGCall.cpp Thu Mar 29 08:02:08 2018
@@ -64,6 +64,7 @@ unsigned CodeGenTypes::ClangCallConvToLL
case CC_PreserveMost: return llvm::CallingConv::PreserveMost;
case CC_PreserveAll: return llvm::CallingConv::PreserveAll;
case CC_Swift: return llvm::CallingConv::Swift;
+ case CC_CUDAKernel: return CGM.getTargetCodeGenInfo().getCUDAKernelCallingConv();
}
}
Modified: cfe/trunk/lib/CodeGen/CGDebugInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGDebugInfo.cpp?rev=328795&r1=328794&r2=328795&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGDebugInfo.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGDebugInfo.cpp Thu Mar 29 08:02:08 2018
@@ -1022,6 +1022,9 @@ static unsigned getDwarfCC(CallingConv C
return llvm::dwarf::DW_CC_LLVM_PreserveAll;
case CC_X86RegCall:
return llvm::dwarf::DW_CC_LLVM_X86RegCall;
+ case CC_CUDAKernel:
+ // ToDo: Add llvm::dwarf::DW_CC_LLVM_CUDAKernel;
+ return 0;
}
return 0;
}
Modified: cfe/trunk/lib/CodeGen/TargetInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/TargetInfo.cpp?rev=328795&r1=328794&r2=328795&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/TargetInfo.cpp (original)
+++ cfe/trunk/lib/CodeGen/TargetInfo.cpp Thu Mar 29 08:02:08 2018
@@ -431,6 +431,10 @@ unsigned TargetCodeGenInfo::getOpenCLKer
return llvm::CallingConv::SPIR_KERNEL;
}
+unsigned TargetCodeGenInfo::getCUDAKernelCallingConv() const {
+ return llvm::CallingConv::C;
+}
+
llvm::Constant *TargetCodeGenInfo::getNullPointer(const CodeGen::CodeGenModule &CGM,
llvm::PointerType *T, QualType QT) const {
return llvm::ConstantPointerNull::get(T);
@@ -7635,6 +7639,7 @@ public:
void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
CodeGen::CodeGenModule &M) const override;
unsigned getOpenCLKernelCallingConv() const override;
+ unsigned getCUDAKernelCallingConv() const override;
llvm::Constant *getNullPointer(const CodeGen::CodeGenModule &CGM,
llvm::PointerType *T, QualType QT) const override;
@@ -7722,6 +7727,10 @@ unsigned AMDGPUTargetCodeGenInfo::getOpe
return llvm::CallingConv::AMDGPU_KERNEL;
}
+unsigned AMDGPUTargetCodeGenInfo::getCUDAKernelCallingConv() const {
+ return llvm::CallingConv::AMDGPU_KERNEL;
+}
+
// Currently LLVM assumes null pointers always have value 0,
// which results in incorrectly transformed IR. Therefore, instead of
// emitting null pointers in private and local address spaces, a null
Modified: cfe/trunk/lib/CodeGen/TargetInfo.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/TargetInfo.h?rev=328795&r1=328794&r2=328795&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/TargetInfo.h (original)
+++ cfe/trunk/lib/CodeGen/TargetInfo.h Thu Mar 29 08:02:08 2018
@@ -223,6 +223,9 @@ public:
/// Get LLVM calling convention for OpenCL kernel.
virtual unsigned getOpenCLKernelCallingConv() const;
+ /// Get LLVM calling convention for CUDA kernel.
+ virtual unsigned getCUDAKernelCallingConv() const;
+
/// Get target specific null pointer.
/// \param T is the LLVM type of the null pointer.
/// \param QT is the clang QualType of the null pointer.
Modified: cfe/trunk/lib/Sema/SemaExpr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExpr.cpp?rev=328795&r1=328794&r2=328795&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaExpr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaExpr.cpp Thu Mar 29 08:02:08 2018
@@ -25,6 +25,7 @@
#include "clang/AST/ExprObjC.h"
#include "clang/AST/ExprOpenMP.h"
#include "clang/AST/RecursiveASTVisitor.h"
+#include "clang/AST/Type.h"
#include "clang/AST/TypeLoc.h"
#include "clang/Basic/PartialDiagnostic.h"
#include "clang/Basic/SourceManager.h"
@@ -1657,6 +1658,16 @@ Sema::BuildDeclRefExpr(ValueDecl *D, Qua
isa<VarDecl>(D) &&
NeedToCaptureVariable(cast<VarDecl>(D), NameInfo.getLoc());
+ // Drop CUDA kernel calling convention since it is invisible to the user
+ // in DRE.
+ if (const auto *FT = Ty->getAs<FunctionType>()) {
+ if (FT->getCallConv() == CC_CUDAKernel) {
+ FT = Context.adjustFunctionType(FT,
+ FT->getExtInfo().withCallingConv(CC_C));
+ Ty = QualType(FT, Ty.getQualifiers().getAsOpaqueValue());
+ }
+ }
+
DeclRefExpr *E;
if (isa<VarTemplateSpecializationDecl>(D)) {
VarTemplateSpecializationDecl *VarSpec =
Modified: cfe/trunk/lib/Sema/SemaOverload.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOverload.cpp?rev=328795&r1=328794&r2=328795&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOverload.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOverload.cpp Thu Mar 29 08:02:08 2018
@@ -1481,7 +1481,6 @@ bool Sema::IsFunctionConversion(QualType
.getTypePtr());
Changed = true;
}
-
// Convert FromFPT's ExtParameterInfo if necessary. The conversion is valid
// only if the ExtParameterInfo lists of the two function prototypes can be
// merged and the merged list is identical to ToFPT's ExtParameterInfo list.
Modified: cfe/trunk/lib/Sema/SemaType.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=328795&r1=328794&r2=328795&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaType.cpp (original)
+++ cfe/trunk/lib/Sema/SemaType.cpp Thu Mar 29 08:02:08 2018
@@ -3316,6 +3316,18 @@ getCCForDeclaratorChunk(Sema &S, Declara
CallingConv CC = S.Context.getDefaultCallingConvention(FTI.isVariadic,
IsCXXInstanceMethod);
+ // Attribute AT_CUDAGlobal affects the calling convention for AMDGPU targets.
+ // This is the simplest place to infer calling convention for CUDA kernels.
+ if (S.getLangOpts().CUDA && S.getLangOpts().CUDAIsDevice) {
+ for (const AttributeList *Attr = D.getDeclSpec().getAttributes().getList();
+ Attr; Attr = Attr->getNext()) {
+ if (Attr->getKind() == AttributeList::AT_CUDAGlobal) {
+ CC = CC_CUDAKernel;
+ break;
+ }
+ }
+ }
+
// Attribute AT_OpenCLKernel affects the calling convention for SPIR
// and AMDGPU targets, hence it cannot be treated as a calling
// convention attribute. This is the simplest place to infer
Added: cfe/trunk/test/CodeGenCUDA/kernel-amdgcn.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/kernel-amdgcn.cu?rev=328795&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/kernel-amdgcn.cu (added)
+++ cfe/trunk/test/CodeGenCUDA/kernel-amdgcn.cu Thu Mar 29 08:02:08 2018
@@ -0,0 +1,29 @@
+// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm %s -o - | FileCheck %s
+#include "Inputs/cuda.h"
+
+// CHECK: define amdgpu_kernel void @_ZN1A6kernelEv
+class A {
+public:
+ static __global__ void kernel(){}
+};
+
+// CHECK: define void @_Z10non_kernelv
+__device__ void non_kernel(){}
+
+// CHECK: define amdgpu_kernel void @_Z6kerneli
+__global__ void kernel(int x) {
+ non_kernel();
+}
+
+// CHECK: define amdgpu_kernel void @_Z15template_kernelI1AEvT_
+template<class T>
+__global__ void template_kernel(T x) {}
+
+void launch(void *f);
+
+int main() {
+ launch((void*)A::kernel);
+ launch((void*)kernel);
+ launch((void*)template_kernel<A>);
+ return 0;
+}
Modified: cfe/trunk/tools/libclang/CXType.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/tools/libclang/CXType.cpp?rev=328795&r1=328794&r2=328795&view=diff
==============================================================================
--- cfe/trunk/tools/libclang/CXType.cpp (original)
+++ cfe/trunk/tools/libclang/CXType.cpp Thu Mar 29 08:02:08 2018
@@ -626,6 +626,7 @@ CXCallingConv clang_getFunctionTypeCalli
TCALLINGCONV(PreserveAll);
case CC_SpirFunction: return CXCallingConv_Unexposed;
case CC_OpenCLKernel: return CXCallingConv_Unexposed;
+ case CC_CUDAKernel: return CXCallingConv_Unexposed;
break;
}
#undef TCALLINGCONV
More information about the cfe-commits
mailing list