From: Peter Collingbourne <peter@pcc.me.uk>
Date: Thu, 6 Oct 2011 16:49:54 +0000 (+0000)
Subject: CUDA: set proper calling conventions for PTX
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=744d90bfe2a43847764a707b1bee7ef1e30ad5f2;p=clang

CUDA: set proper calling conventions for PTX

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@141296 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/lib/CodeGen/TargetInfo.cpp b/lib/CodeGen/TargetInfo.cpp
index 83e01132de..91802d3b3a 100644
--- a/lib/CodeGen/TargetInfo.cpp
+++ b/lib/CodeGen/TargetInfo.cpp
@@ -2774,8 +2774,9 @@ void PTXABIInfo::computeInfo(CGFunctionInfo &FI) const {
 
   // 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 @@ void PTXTargetCodeGenInfo::SetTargetAttributes(const Decl *D,
   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);
   }
 }
 
diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
index e68d0cf6c3..b7356c2de0 100644
--- a/test/CMakeLists.txt
+++ b/test/CMakeLists.txt
@@ -2,6 +2,7 @@ set(CLANG_TEST_DIRECTORIES
   "Analysis"
   "CodeCompletion"
   "CodeGen"
+  "CodeGenCUDA"
   "CodeGenCXX"
   "CodeGenObjC"
   "CodeGenOpenCL"
diff --git a/test/CodeGenCUDA/ptx-kernels.cu b/test/CodeGenCUDA/ptx-kernels.cu
new file mode 100644
index 0000000000..310fa2a276
--- /dev/null
+++ b/test/CodeGenCUDA/ptx-kernels.cu
@@ -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();
+}