[clang] e3e47e8 - [OpenMP] Make complex soft-float functions on the GPU weak definitions

[Johannes Doerfert via cfe-commits]

Author: Johannes Doerfert
Date: 2020-07-09T01:06:55-05:00
New Revision: e3e47e80355422df2e730cf97a0c80bb6de3915e

URL: https://github.com/llvm/llvm-project/commit/e3e47e80355422df2e730cf97a0c80bb6de3915e
DIFF: https://github.com/llvm/llvm-project/commit/e3e47e80355422df2e730cf97a0c80bb6de3915e.diff

LOG: [OpenMP] Make complex soft-float functions on the GPU weak definitions

To avoid linkage errors we have to ensure the linkage allows multiple
definitions of these compiler inserted functions. Since they are on the
cold path of complex computations, we want to avoid `inline`. Instead,
we opt for `weak` and `noinline` for now.

Added: 
    

Modified: 
    clang/lib/Headers/__clang_cuda_complex_builtins.h
    clang/test/Headers/nvptx_device_math_complex.c
    clang/test/Headers/nvptx_device_math_complex.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/Headers/__clang_cuda_complex_builtins.h b/clang/lib/Headers/__clang_cuda_complex_builtins.h
index d698be71d011..c48c754ed1a4 100644
--- a/clang/lib/Headers/__clang_cuda_complex_builtins.h
+++ b/clang/lib/Headers/__clang_cuda_complex_builtins.h
@@ -18,7 +18,7 @@
 #pragma push_macro("__DEVICE__")
 #ifdef _OPENMP
 #pragma omp declare target
-#define __DEVICE__ __attribute__((noinline, nothrow, cold))
+#define __DEVICE__ __attribute__((noinline, nothrow, cold, weak))
 #else
 #define __DEVICE__ __device__ inline
 #endif

diff  --git a/clang/test/Headers/nvptx_device_math_complex.c b/clang/test/Headers/nvptx_device_math_complex.c
index 9b96b5dd8c22..0e212592dd2b 100644
--- a/clang/test/Headers/nvptx_device_math_complex.c
+++ b/clang/test/Headers/nvptx_device_math_complex.c
@@ -11,10 +11,10 @@
 #include <complex.h>
 #endif
 
-// CHECK-DAG: define {{.*}} @__mulsc3
-// CHECK-DAG: define {{.*}} @__muldc3
-// CHECK-DAG: define {{.*}} @__divsc3
-// CHECK-DAG: define {{.*}} @__divdc3
+// CHECK-DAG: define weak {{.*}} @__mulsc3
+// CHECK-DAG: define weak {{.*}} @__muldc3
+// CHECK-DAG: define weak {{.*}} @__divsc3
+// CHECK-DAG: define weak {{.*}} @__divdc3
 
 // CHECK-DAG: call float @__nv_scalbnf(
 void test_scmplx(float _Complex a) {

diff  --git a/clang/test/Headers/nvptx_device_math_complex.cpp b/clang/test/Headers/nvptx_device_math_complex.cpp
index 15434d907605..58ed24b74b0e 100644
--- a/clang/test/Headers/nvptx_device_math_complex.cpp
+++ b/clang/test/Headers/nvptx_device_math_complex.cpp
@@ -5,10 +5,10 @@
 
 #include <complex>
 
-// CHECK-DAG: define {{.*}} @__mulsc3
-// CHECK-DAG: define {{.*}} @__muldc3
-// CHECK-DAG: define {{.*}} @__divsc3
-// CHECK-DAG: define {{.*}} @__divdc3
+// CHECK-DAG: define weak {{.*}} @__mulsc3
+// CHECK-DAG: define weak {{.*}} @__muldc3
+// CHECK-DAG: define weak {{.*}} @__divsc3
+// CHECK-DAG: define weak {{.*}} @__divdc3
 
 // CHECK-DAG: call float @__nv_scalbnf(
 void test_scmplx(std::complex<float> a) {