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
diff --git a/lib/CodeGen/TargetInfo.cpp b/lib/CodeGen/TargetInfo.cpp
index 9debf5e..83e0113 100644
--- a/lib/CodeGen/TargetInfo.cpp
+++ b/lib/CodeGen/TargetInfo.cpp
@@ -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);
+ }
+}
+
}
//===----------------------------------------------------------------------===//
diff --git a/test/CodeGenOpenCL/ptx-calls.cl b/test/CodeGenOpenCL/ptx-calls.cl
new file mode 100644
index 0000000..6f33640
--- /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 0000000..4d6fa10
--- /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()
+
