maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[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.
This commit is contained in:
Johannes Doerfert 2020-07-09 01:04:16 -05:00
parent 305b500eaf
commit e3e47e8035
3 changed files with 9 additions and 9 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
clang/lib/Headers/__clang_cuda_complex_builtins.h
View File

@ -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

8
clang/test/Headers/nvptx_device_math_complex.c
View File

@ -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) {

8
clang/test/Headers/nvptx_device_math_complex.cpp
View File

@ -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) {