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

3 files changed, 62 insertions, 5 deletions

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<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 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()
+