From 818eafb6ac56c87b80b34be29ca115cd309026d2 Mon Sep 17 00:00:00 2001 From: Justin Holewinski Date: Wed, 5 Oct 2011 17:58:44 +0000 Subject: [PATCH] PTX: Set proper calling conventions for PTX in OpenCL mode. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@141193 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/CodeGen/TargetInfo.cpp | 45 +++++++++++++++++++++++++++---- test/CodeGenOpenCL/ptx-calls.cl | 12 +++++++++ test/CodeGenOpenCL/ptx-kernels.cl | 10 +++++++ 3 files changed, 62 insertions(+), 5 deletions(-) create mode 100644 test/CodeGenOpenCL/ptx-calls.cl create mode 100644 test/CodeGenOpenCL/ptx-kernels.cl diff --git a/lib/CodeGen/TargetInfo.cpp b/lib/CodeGen/TargetInfo.cpp index 9debf5ef4d..83e01132de 100644 --- a/lib/CodeGen/TargetInfo.cpp +++ b/lib/CodeGen/TargetInfo.cpp @@ -2742,6 +2742,9 @@ class PTXTargetCodeGenInfo : public TargetCodeGenInfo { public: PTXTargetCodeGenInfo(CodeGenTypes &CGT) : TargetCodeGenInfo(new PTXABIInfo(CGT)) {} + + virtual void SetTargetAttributes(const Decl *D, llvm::GlobalValue *GV, + CodeGen::CodeGenModule &M) const; }; ABIArgInfo PTXABIInfo::classifyReturnType(QualType RetTy) const { @@ -2771,13 +2774,20 @@ void PTXABIInfo::computeInfo(CGFunctionInfo &FI) const { // Calling convention as default by an ABI. llvm::CallingConv::ID DefaultCC; - StringRef Env = getContext().getTargetInfo().getTriple().getEnvironmentName(); - if (Env == "device") + if (getContext().getLangOptions().OpenCL) { + // If we are in OpenCL mode, then default to device functions DefaultCC = llvm::CallingConv::PTX_Device; - else - DefaultCC = llvm::CallingConv::PTX_Kernel; - + } else { + // If we are in standard C/C++ mode, use the triple to decide on the default + StringRef Env = + getContext().getTargetInfo().getTriple().getEnvironmentName(); + if (Env == "device") + DefaultCC = llvm::CallingConv::PTX_Device; + else + DefaultCC = llvm::CallingConv::PTX_Kernel; + } FI.setEffectiveCallingConvention(DefaultCC); + } llvm::Value *PTXABIInfo::EmitVAArg(llvm::Value *VAListAddr, QualType Ty, @@ -2786,6 +2796,31 @@ llvm::Value *PTXABIInfo::EmitVAArg(llvm::Value *VAListAddr, QualType Ty, return 0; } +void PTXTargetCodeGenInfo::SetTargetAttributes(const Decl *D, + llvm::GlobalValue *GV, + CodeGen::CodeGenModule &M) const{ + const FunctionDecl *FD = dyn_cast(D); + if (!FD) return; + + llvm::Function *F = cast(GV); + + // Perform special handling in OpenCL mode + if (M.getContext().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; + // And kernel functions are not subject to inlining + F->addFnAttr(llvm::Attribute::NoInline); + } + + // Set the derived calling convention + F->setCallingConv(CC); + } +} + } //===----------------------------------------------------------------------===// diff --git a/test/CodeGenOpenCL/ptx-calls.cl b/test/CodeGenOpenCL/ptx-calls.cl new file mode 100644 index 0000000000..6f336405c3 --- /dev/null +++ b/test/CodeGenOpenCL/ptx-calls.cl @@ -0,0 +1,12 @@ +// RUN: %clang_cc1 %s -triple ptx32-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s + +void device_function() { +} +// CHECK: define ptx_device void @device_function() + +__kernel void kernel_function() { + device_function(); +} +// CHECK: define ptx_kernel void @kernel_function() +// CHECK: call ptx_device void @device_function() + diff --git a/test/CodeGenOpenCL/ptx-kernels.cl b/test/CodeGenOpenCL/ptx-kernels.cl new file mode 100644 index 0000000000..4d6fa1084d --- /dev/null +++ b/test/CodeGenOpenCL/ptx-kernels.cl @@ -0,0 +1,10 @@ +// RUN: %clang_cc1 %s -triple ptx32-unknown-unknown -emit-llvm -o - | FileCheck %s + +void device_function() { +} +// CHECK: define ptx_device void @device_function() + +__kernel void kernel_function() { +} +// CHECK: define ptx_kernel void @kernel_function() + -- 2.40.0