diff options
author | Artem Belevich <tra@google.com> | 2015-10-27 17:56:59 +0000 |
---|---|---|
committer | Artem Belevich <tra@google.com> | 2015-10-27 17:56:59 +0000 |
commit | d40c9d5c5ac1cfe4493652f3d0e3bfbded161840 (patch) | |
tree | 51d349f5fb4bf134849e4ca8de0d16ff4eff4a09 /test | |
parent | bbf167646bc113cbf33bc9e6e5bdf99314187e49 (diff) | |
download | clang-d40c9d5c5ac1cfe4493652f3d0e3bfbded161840.tar.gz |
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
Diffstat (limited to 'test')
-rw-r--r-- | test/CodeGen/link-bitcode-file.c | 15 | ||||
-rw-r--r-- | test/CodeGenCUDA/Inputs/device-code-2.ll | 16 | ||||
-rw-r--r-- | test/CodeGenCUDA/link-device-bitcode.cu | 18 |
3 files changed, 45 insertions, 4 deletions
diff --git a/test/CodeGen/link-bitcode-file.c b/test/CodeGen/link-bitcode-file.c index 92b1a88ffb..7810fe1d29 100644 --- a/test/CodeGen/link-bitcode-file.c +++ b/test/CodeGen/link-bitcode-file.c @@ -1,6 +1,12 @@ // RUN: %clang_cc1 -triple i386-pc-linux-gnu -DBITCODE -emit-llvm-bc -o %t.bc %s -// RUN: %clang_cc1 -triple i386-pc-linux-gnu -mlink-bitcode-file %t.bc -O3 -emit-llvm -o - %s | FileCheck -check-prefix=CHECK-NO-BC %s -// RUN: not %clang_cc1 -triple i386-pc-linux-gnu -DBITCODE -mlink-bitcode-file %t.bc -O3 -emit-llvm -o - %s 2>&1 | FileCheck -check-prefix=CHECK-BC %s +// RUN: %clang_cc1 -triple i386-pc-linux-gnu -DBITCODE2 -emit-llvm-bc -o %t-2.bc %s +// RUN: %clang_cc1 -triple i386-pc-linux-gnu -mlink-bitcode-file %t.bc \ +// RUN: -O3 -emit-llvm -o - %s | FileCheck -check-prefix=CHECK-NO-BC %s +// RUN: %clang_cc1 -triple i386-pc-linux-gnu -O3 -emit-llvm -o - \ +// RUN: -mlink-bitcode-file %t.bc -mlink-bitcode-file %t-2.bc %s \ +// RUN: | FileCheck -check-prefix=CHECK-NO-BC -check-prefix=CHECK-NO-BC2 %s +// RUN: not %clang_cc1 -triple i386-pc-linux-gnu -DBITCODE -O3 -emit-llvm -o - \ +// RUN: -mlink-bitcode-file %t.bc %s 2>&1 | FileCheck -check-prefix=CHECK-BC %s // Make sure we deal with failure to load the file. // RUN: not %clang_cc1 -triple i386-pc-linux-gnu -mlink-bitcode-file no-such-file.bc \ // RUN: -emit-llvm -o - %s 2>&1 | FileCheck -check-prefix=CHECK-NO-FILE %s @@ -9,11 +15,15 @@ int f(void); #ifdef BITCODE +extern int f2(void); // CHECK-BC: fatal error: cannot link module {{.*}}'f': symbol multiply defined int f(void) { + f2(); return 42; } +#elif BITCODE2 +int f2(void) { return 43; } #else // CHECK-NO-BC-LABEL: define i32 @g @@ -23,6 +33,7 @@ int g(void) { } // CHECK-NO-BC-LABEL: define i32 @f +// CHECK-NO-BC2-LABEL: define i32 @f2 #endif 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 |