[cfe-commits] r141193 - in /cfe/trunk: lib/CodeGen/TargetInfo.cpp test/CodeGenOpenCL/ptx-calls.cl test/CodeGenOpenCL/ptx-kernels.cl
Justin Holewinski
justin.holewinski at gmail.com
Wed Oct 5 10:58:47 PDT 2011
Author: jholewinski
Date: Wed Oct 5 12:58:44 2011
New Revision: 141193
URL: http://llvm.org/viewvc/llvm-project?rev=141193&view=rev
Log:
PTX: Set proper calling conventions for PTX in OpenCL mode.
Added:
cfe/trunk/test/CodeGenOpenCL/ptx-calls.cl
cfe/trunk/test/CodeGenOpenCL/ptx-kernels.cl
Modified:
cfe/trunk/lib/CodeGen/TargetInfo.cpp
Modified: cfe/trunk/lib/CodeGen/TargetInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/TargetInfo.cpp?rev=141193&r1=141192&r2=141193&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/TargetInfo.cpp (original)
+++ cfe/trunk/lib/CodeGen/TargetInfo.cpp Wed Oct 5 12:58:44 2011
@@ -2742,6 +2742,9 @@
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 @@
// 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 @@
return 0;
}
+void PTXTargetCodeGenInfo::SetTargetAttributes(const Decl *D,
+ llvm::GlobalValue *GV,
+ CodeGen::CodeGenModule &M) const{
+ const FunctionDecl *FD = dyn_cast<FunctionDecl>(D);
+ if (!FD) return;
+
+ llvm::Function *F = cast<llvm::Function>(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<OpenCLKernelAttr>()) {
+ // 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);
+ }
+}
+
}
//===----------------------------------------------------------------------===//
Added: cfe/trunk/test/CodeGenOpenCL/ptx-calls.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/ptx-calls.cl?rev=141193&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/ptx-calls.cl (added)
+++ cfe/trunk/test/CodeGenOpenCL/ptx-calls.cl Wed Oct 5 12:58:44 2011
@@ -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()
+
Added: cfe/trunk/test/CodeGenOpenCL/ptx-kernels.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/ptx-kernels.cl?rev=141193&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/ptx-kernels.cl (added)
+++ cfe/trunk/test/CodeGenOpenCL/ptx-kernels.cl Wed Oct 5 12:58:44 2011
@@ -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()
+
More information about the cfe-commits
mailing list