From 5bad4afa2fd86a90c42d92cc7a24bec269365c8b Mon Sep 17 00:00:00 2001 From: Peter Collingbourne Date: Thu, 6 Oct 2011 16:49:54 +0000 Subject: CUDA: set proper calling conventions for PTX llvm-svn: 141296 --- clang/lib/CodeGen/TargetInfo.cpp | 20 +++++++++++++------- 1 file changed, 13 insertions(+), 7 deletions(-) (limited to 'clang/lib/CodeGen/TargetInfo.cpp') diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp index 83e01132dea..91802d3b3ae 100644 --- a/clang/lib/CodeGen/TargetInfo.cpp +++ b/clang/lib/CodeGen/TargetInfo.cpp @@ -2774,8 +2774,9 @@ void PTXABIInfo::computeInfo(CGFunctionInfo &FI) const { // Calling convention as default by an ABI. llvm::CallingConv::ID DefaultCC; - if (getContext().getLangOptions().OpenCL) { - // If we are in OpenCL mode, then default to device functions + const LangOptions &LangOpts = getContext().getLangOptions(); + if (LangOpts.OpenCL || LangOpts.CUDA) { + // If we are in OpenCL or CUDA mode, then default to device functions DefaultCC = llvm::CallingConv::PTX_Device; } else { // If we are in standard C/C++ mode, use the triple to decide on the default @@ -2805,19 +2806,24 @@ void PTXTargetCodeGenInfo::SetTargetAttributes(const Decl *D, llvm::Function *F = cast(GV); // Perform special handling in OpenCL mode - if (M.getContext().getLangOptions().OpenCL) { + if (M.getLangOptions().OpenCL) { // Use OpenCL function attributes to set proper calling conventions // By default, all functions are device functions - llvm::CallingConv::ID CC = llvm::CallingConv::PTX_Device; if (FD->hasAttr()) { // OpenCL __kernel functions get a kernel calling convention - CC = llvm::CallingConv::PTX_Kernel; + F->setCallingConv(llvm::CallingConv::PTX_Kernel); // And kernel functions are not subject to inlining F->addFnAttr(llvm::Attribute::NoInline); } + } - // Set the derived calling convention - F->setCallingConv(CC); + // Perform special handling in CUDA mode. + if (M.getLangOptions().CUDA) { + // CUDA __global__ functions get a kernel calling convention. Since + // __global__ functions cannot be called from the device, we do not + // need to set the noinline attribute. + if (FD->getAttr()) + F->setCallingConv(llvm::CallingConv::PTX_Kernel); } } -- cgit v1.2.3