[cfe-commits] r141296 - in /cfe/trunk: lib/CodeGen/TargetInfo.cpp test/CMakeLists.txt test/CodeGenCUDA/ test/CodeGenCUDA/ptx-kernels.cu

[Peter Collingbourne]

Author: pcc
Date: Thu Oct  6 11:49:54 2011
New Revision: 141296

URL: http://llvm.org/viewvc/llvm-project?rev=141296&view=rev
Log:
CUDA: set proper calling conventions for PTX

Added:
    cfe/trunk/test/CodeGenCUDA/
    cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
Modified:
    cfe/trunk/lib/CodeGen/TargetInfo.cpp
    cfe/trunk/test/CMakeLists.txt

Modified: cfe/trunk/lib/CodeGen/TargetInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/TargetInfo.cpp?rev=141296&r1=141295&r2=141296&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/TargetInfo.cpp (original)
+++ cfe/trunk/lib/CodeGen/TargetInfo.cpp Thu Oct  6 11:49:54 2011
@@ -2774,8 +2774,9 @@
 
   // Calling convention as default by an ABI.
   llvm::CallingConv::ID DefaultCC;
-  if (getContext().getLangOptions().OpenCL) {
-    // If we are in OpenCL mode, then default to device functions
+  const LangOptions &LangOpts = getContext().getLangOptions();
+  if (LangOpts.OpenCL || LangOpts.CUDA) {
+    // If we are in OpenCL or CUDA mode, then default to device functions
     DefaultCC = llvm::CallingConv::PTX_Device;
   } else {
     // If we are in standard C/C++ mode, use the triple to decide on the default
@@ -2805,19 +2806,24 @@
   llvm::Function *F = cast<llvm::Function>(GV);
 
   // Perform special handling in OpenCL mode
-  if (M.getContext().getLangOptions().OpenCL) {
+  if (M.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;
+      F->setCallingConv(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);
+  // Perform special handling in CUDA mode.
+  if (M.getLangOptions().CUDA) {
+    // CUDA __global__ functions get a kernel calling convention.  Since
+    // __global__ functions cannot be called from the device, we do not
+    // need to set the noinline attribute.
+    if (FD->getAttr<CUDAGlobalAttr>())
+      F->setCallingConv(llvm::CallingConv::PTX_Kernel);
   }
 }
 

Modified: cfe/trunk/test/CMakeLists.txt
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CMakeLists.txt?rev=141296&r1=141295&r2=141296&view=diff
==============================================================================
--- cfe/trunk/test/CMakeLists.txt (original)
+++ cfe/trunk/test/CMakeLists.txt Thu Oct  6 11:49:54 2011
@@ -2,6 +2,7 @@
   "Analysis"
   "CodeCompletion"
   "CodeGen"
+  "CodeGenCUDA"
   "CodeGenCXX"
   "CodeGenObjC"
   "CodeGenOpenCL"

Added: cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu?rev=141296&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu (added)
+++ cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu Thu Oct  6 11:49:54 2011
@@ -0,0 +1,12 @@
+// RUN: %clang_cc1 %s -triple ptx32-unknown-unknown -emit-llvm -o - | FileCheck %s
+
+#include "../SemaCUDA/cuda.h"
+
+// CHECK: define ptx_device{{.*}}device_function
+__device__ void device_function() {}
+
+// CHECK: define ptx_kernel{{.*}}global_function
+__global__ void global_function() {
+  // CHECK: call ptx_device{{.*}}device_function
+  device_function();
+}