[cfe-commits] r141193 - in /cfe/trunk: lib/CodeGen/TargetInfo.cpp test/CodeGenOpenCL/ptx-calls.cl test/CodeGenOpenCL/ptx-kernels.cl

[Justin Holewinski]

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