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 83e0113..91802d3 100644
--- a/lib/CodeGen/TargetInfo.cpp
+++ b/lib/CodeGen/TargetInfo.cpp
@@ -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);
}
}
diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
index e68d0cf..b7356c2 100644
--- a/test/CMakeLists.txt
+++ b/test/CMakeLists.txt
@@ -2,6 +2,7 @@
"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 0000000..310fa2a
--- /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();
+}