Allow linking multiple bitcode files.

Linking options for particular file depend on the option that specifies the file.
Currently there are two:

* -mlink-bitcode-file links in complete content of the specified file.
* -mlink-cuda-bitcode links in only the symbols needed by current TU.
   Linked symbols are internalized. This bitcode linking mode is used to
   link device-specific bitcode provided by CUDA.

Files are linked in order they are specified on command line.

-mlink-cuda-bitcode replaces -fcuda-uses-libdevice flag.

Differential Revision: http://reviews.llvm.org/D13913

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

2 files changed, 32 insertions, 2 deletions

diff --git a/test/CodeGenCUDA/Inputs/device-code-2.ll b/test/CodeGenCUDA/Inputs/device-code-2.ll
new file mode 100644
index 0000000000..8fde3b13ec
--- /dev/null
+++ b/test/CodeGenCUDA/Inputs/device-code-2.ll
@@ -0,0 +1,16 @@
+; Simple bit of IR to mimic CUDA's libdevice.
+
+target triple = "nvptx-unknown-cuda"
+
+define double @__nv_sin(double %a) {
+       ret double 1.0
+}
+
+define double @__nv_exp(double %a) {
+       ret double 3.0
+}
+
+define double @__unused(double %a) {
+       ret double 2.0
+}
+
diff --git a/test/CodeGenCUDA/link-device-bitcode.cu b/test/CodeGenCUDA/link-device-bitcode.cu
index 45e5bcff99..de3d39c20b 100644
--- a/test/CodeGenCUDA/link-device-bitcode.cu
+++ b/test/CodeGenCUDA/link-device-bitcode.cu
@@ -6,13 +6,21 @@
 // Prepare bitcode file to link with
 // RUN: %clang_cc1 -triple nvptx-unknown-cuda -emit-llvm-bc -o %t.bc \
 // RUN:    %S/Inputs/device-code.ll
+// RUN: %clang_cc1 -triple nvptx-unknown-cuda -emit-llvm-bc -o %t-2.bc \
+// RUN:    %S/Inputs/device-code-2.ll
 //
 // Make sure function in device-code gets linked in and internalized.
 // RUN: %clang_cc1 -triple nvptx-unknown-cuda -fcuda-is-device \
-// RUN:    -mlink-bitcode-file %t.bc -fcuda-uses-libdevice -emit-llvm \
+// RUN:    -mlink-cuda-bitcode %t.bc  -emit-llvm \
 // RUN:    -disable-llvm-passes -o - %s \
 // RUN:    | FileCheck %s -check-prefix CHECK-IR
 //
+// Make sure we can link two bitcode files.
+// RUN: %clang_cc1 -triple nvptx-unknown-cuda -fcuda-is-device \
+// RUN:    -mlink-cuda-bitcode %t.bc -mlink-cuda-bitcode %t-2.bc \
+// RUN:    -emit-llvm -disable-llvm-passes -o - %s \
+// RUN:    | FileCheck %s -check-prefix CHECK-IR -check-prefix CHECK-IR-2
+//
 // Make sure function in device-code gets linked but is not internalized
 // without -fcuda-uses-libdevice
 // RUN: %clang_cc1 -triple nvptx-unknown-cuda -fcuda-is-device \
@@ -22,7 +30,7 @@
 //
 // Make sure NVVMReflect pass is enabled in NVPTX back-end.
 // RUN: %clang_cc1 -triple nvptx-unknown-cuda -fcuda-is-device \
-// RUN:    -mlink-bitcode-file %t.bc -fcuda-uses-libdevice -S -o /dev/null %s \
+// RUN:    -mlink-cuda-bitcode %t.bc -S -o /dev/null %s \
 // RUN:    -backend-option -debug-pass=Structure 2>&1 \
 // RUN:    | FileCheck %s -check-prefix CHECK-REFLECT
 
@@ -52,5 +60,11 @@ __global__ __attribute__((used)) void kernel(float *out, float *in) {
 // CHECK-IR: call i32 @__nvvm_reflect
 // CHECK-IR: ret float
 
+// Make sure we've linked in and internalized only needed functions
+// from the second bitcode file.
+// CHECK-IR-2-LABEL: define internal double @__nv_sin
+// CHECK-IR-2-LABEL: define internal double @__nv_exp
+// CHECK-IR-2-NOT: double @__unused
+
 // Verify that NVVMReflect pass is among the passes run by NVPTX back-end.
 // CHECK-REFLECT: Replace occurrences of __nvvm_reflect() calls with 0/1