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) { _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
