[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.

diff --git a/clang/lib/Headers/__clang_cuda_complex_builtins.h b/clang/lib/Headers/__clang_cuda_complex_builtins.h
index d698be7..c48c754 100644 (file)
--- 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 9b96b5d..0e21259 100644 (file)
--- 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 15434d9..58ed24b 100644 (file)
--- 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) {