fsanitize flag fails on amdgpu

[Quuxplusone]

Bugzilla Link	PR46422
Status	NEW
Importance	P normal
Reported by	Ahmad Lashgar (ahmad.lashgar@barco.com)
Reported on	2020-06-22 07:38:25 -0700
Last modified on	2020-10-21 11:50:09 -0700
Version	trunk
Hardware	PC Linux
CC	llvm-bugs@lists.llvm.org, tra@google.com
Fixed by commit(s)
Attachments
Blocks
Blocked by
See also

Trying to compile simple vector-add GPU kernel for amdgpu with -
fsanitize=thread flag. Linker fails to link against libtsan with a few
undefined reference errors (this is with llvm trunk build based on source from
April 29, 2020). I verified that compiling with fsanitize works for none-GPU
code. Is fsanitize supported on amdgpu?

Here is the full source code:

////////////////
#include <hip/hip_runtime.h>

__global__ void vector_add(float *out, float *a, float *b, int n)
{
  int i = threadIdx.x + blockDim.x * blockIdx.x;
  if (i < n)
  {
    out[i] = a[i] + b[i];
  }
}
////////////////

I passed -v to compiler to get more details on steps and here is the compiler
command and full verbose log:

"/home/ahmlas/opt/llvm-200429-release/bin/clang++" -cc1 -triple amdgcn-amd-
amdhsa -aux-triple x86_64-unknown-linux-gnu -emit-llvm-bc -emit-llvm-uselists -
disable-free -disable-llvm-verifier -discard-value-names -main-file-name vector-
add.cu -mrelocation-model pic -pic-level 1 -mthread-model posix -mframe-
pointer=none -fno-rounding-math -mconstructor-aliases -munwind-tables -aux-
target-cpu x86-64 -target-cpu gfx900 -fcuda-is-device -fcuda-allow-variadic-
functions -fvisibility hidden -fapply-global-visibility-to-externs -mlink-
builtin-bitcode /opt/rocm-3.3.0/lib/hip.amdgcn.bc -mlink-builtin-bitcode
/opt/rocm-3.3.0/lib/ocml.amdgcn.bc -mlink-builtin-bitcode /opt/rocm-
3.3.0/lib/ockl.amdgcn.bc -mlink-builtin-bitcode /opt/rocm-
3.3.0/lib/oclc_finite_only_off.amdgcn.bc -mlink-builtin-bitcode /opt/rocm-
3.3.0/lib/oclc_daz_opt_off.amdgcn.bc -mlink-builtin-bitcode /opt/rocm-
3.3.0/lib/oclc_correctly_rounded_sqrt_on.amdgcn.bc -mlink-builtin-bitcode
/opt/rocm-3.3.0/lib/oclc_unsafe_math_off.amdgcn.bc -mlink-builtin-bitcode
/opt/rocm-3.3.0/lib/oclc_isa_version_900.amdgcn.bc -mlink-builtin-bitcode
/opt/rocm-3.3.0/lib/oclc_wavefrontsize64_on.amdgcn.bc -dwarf-column-info -fno-
split-dwarf-inlining -debugger-tuning=gdb -v -resource-dir
/home/ahmlas/opt/llvm-200429-release/lib/clang/11.0.0 -isystem
/home/ahmlas/opt/llvm-200429-release/lib/clang/11.0.0/include -isystem
/opt/rocm-3.3.0/hsa/include -isystem /opt/rocm-3.3.0/hip/include -D
__HIP_ARCH_GFX900__=1 -internal-isystem /usr/lib/gcc/x86_64-linux-
gnu/7.5.0/../../../../include/c++/7.5.0 -internal-isystem /usr/lib/gcc/x86_64-
linux-gnu/7.5.0/../../../../include/x86_64-linux-gnu/c++/7.5.0 -internal-
isystem /usr/lib/gcc/x86_64-linux-gnu/7.5.0/../../../../include/x86_64-linux-
gnu/c++/7.5.0 -internal-isystem /usr/lib/gcc/x86_64-linux-
gnu/7.5.0/../../../../include/c++/7.5.0/backward -internal-isystem
/usr/lib/gcc/x86_64-linux-gnu/7.5.0/../../../../include/c++/7.5.0 -internal-
isystem /usr/lib/gcc/x86_64-linux-gnu/7.5.0/../../../../include/x86_64-linux-
gnu/c++/7.5.0 -internal-isystem /usr/lib/gcc/x86_64-linux-
gnu/7.5.0/../../../../include/x86_64-linux-gnu/c++/7.5.0 -internal-isystem
/usr/lib/gcc/x86_64-linux-gnu/7.5.0/../../../../include/c++/7.5.0/backward -
internal-isystem /usr/local/include -internal-isystem /home/ahmlas/opt/llvm-
200429-release/lib/clang/11.0.0/include -internal-externc-isystem
/usr/include/x86_64-linux-gnu -internal-externc-isystem /include -internal-
externc-isystem /usr/include -internal-isystem /usr/local/include -internal-
isystem /home/ahmlas/opt/llvm-200429-release/lib/clang/11.0.0/include -internal-
externc-isystem /usr/include/x86_64-linux-gnu -internal-externc-isystem
/include -internal-externc-isystem /usr/include -O3 -std=c++14 -fdeprecated-
macro -fno-autolink -fdebug-compilation-dir /home/ahmlas/work/src/hipcc-sample -
ferror-limit 19 -fsanitize=thread -fgnuc-version=4.2.1 -fcxx-exceptions -
fexceptions -vectorize-loops -vectorize-slp -mllvm -amdgpu-early-inline-
all=true -mllvm -amdgpu-function-calls=false -fcuda-allow-variadic-functions -
faddrsig -o /tmp/vector-add-02a7b8.bc -x hip vector-add.cu
clang -cc1 version 11.0.0 based upon LLVM 11.0.0git default target x86_64-
unknown-linux-gnu
ignoring nonexistent directory "/include"
ignoring nonexistent directory "/include"
ignoring duplicate directory "/usr/lib/gcc/x86_64-linux-
gnu/7.5.0/../../../../include/x86_64-linux-gnu/c++/7.5.0"
ignoring duplicate directory "/usr/lib/gcc/x86_64-linux-
gnu/7.5.0/../../../../include/c++/7.5.0"
ignoring duplicate directory "/usr/lib/gcc/x86_64-linux-
gnu/7.5.0/../../../../include/x86_64-linux-gnu/c++/7.5.0"
ignoring duplicate directory "/usr/lib/gcc/x86_64-linux-
gnu/7.5.0/../../../../include/x86_64-linux-gnu/c++/7.5.0"
ignoring duplicate directory "/usr/lib/gcc/x86_64-linux-
gnu/7.5.0/../../../../include/c++/7.5.0/backward"
ignoring duplicate directory "/home/ahmlas/opt/llvm-200429-
release/lib/clang/11.0.0/include"
ignoring duplicate directory "/usr/local/include"
ignoring duplicate directory "/home/ahmlas/opt/llvm-200429-
release/lib/clang/11.0.0/include"
ignoring duplicate directory "/usr/include/x86_64-linux-gnu"
ignoring duplicate directory "/usr/include"
ignoring duplicate directory "/usr/local/include"
ignoring duplicate directory "/home/ahmlas/opt/llvm-200429-
release/lib/clang/11.0.0/include"
ignoring duplicate directory "/usr/include"
#include "..." search starts here:
#include <...> search starts here:
 /home/ahmlas/opt/llvm-200429-release/lib/clang/11.0.0/include
 /opt/rocm-3.3.0/hsa/include
 /opt/rocm-3.3.0/hip/include
 /usr/lib/gcc/x86_64-linux-gnu/7.5.0/../../../../include/c++/7.5.0
 /usr/lib/gcc/x86_64-linux-gnu/7.5.0/../../../../include/x86_64-linux-gnu/c++/7.5.0
 /usr/lib/gcc/x86_64-linux-gnu/7.5.0/../../../../include/c++/7.5.0/backward
 /usr/local/include
 /usr/include/x86_64-linux-gnu
 /usr/include
End of search list.
'+fp64-fp16-denormals' is not a recognized feature for this target (ignoring
feature)
'-fp32-denormals' is not a recognized feature for this target (ignoring feature)
'+fp64-fp16-denormals' is not a recognized feature for this target (ignoring
feature)
'-fp32-denormals' is not a recognized feature for this target (ignoring feature)
 "/home/ahmlas/opt/llvm-200429-release/bin/llvm-link" /tmp/vector-add-02a7b8.bc -o /tmp/vector-add-02a7b8-gfx900-linked-117ed1.bc
 "/home/ahmlas/opt/llvm-200429-release/bin/opt" /tmp/vector-add-02a7b8-gfx900-linked-117ed1.bc -O3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -amdgpu-early-inline-all=true -amdgpu-function-calls=false -o /tmp/vector-add-02a7b8-gfx900-optimized-a23c07.bc
 "/home/ahmlas/opt/llvm-200429-release/bin/llc" /tmp/vector-add-02a7b8-gfx900-optimized-a23c07.bc -O3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -filetype=obj -amdgpu-early-inline-all=true -amdgpu-function-calls=false -o /tmp/vector-add-02a7b8-gfx900-d573d4.o
 "/home/ahmlas/opt/llvm-200429-release/bin/lld" -flavor gnu --no-undefined -shared -o /tmp/vector-add-ed90fa.out /tmp/vector-add-02a7b8-gfx900-d573d4.o
lld: error: undefined symbol: __tsan_init
>>> referenced by /tmp/vector-add-02a7b8-gfx900-d573d4.o:(tsan.module_ctor)
>>> referenced by /tmp/vector-add-02a7b8-gfx900-d573d4.o:(tsan.module_ctor)

lld: error: undefined symbol: __tsan_func_entry
>>> referenced by /tmp/vector-add-02a7b8-gfx900-d573d4.o:(vector_add(float*,
float*, float*, int))
>>> referenced by /tmp/vector-add-02a7b8-gfx900-d573d4.o:(vector_add(float*,
float*, float*, int))

lld: error: undefined symbol: __tsan_func_exit
>>> referenced by /tmp/vector-add-02a7b8-gfx900-d573d4.o:(vector_add(float*,
float*, float*, int))
>>> referenced by /tmp/vector-add-02a7b8-gfx900-d573d4.o:(vector_add(float*,
float*, float*, int))
clang++: error: amdgcn-link command failed with exit code 1 (use -v to see
invocation)
Makefile:3: recipe for target 'vector-add.o' failed
make: *** [vector-add.o] Error 1

[Quuxplusone]

This may have been fixed already. At least I can't reproduce it with the clang built from recent sources.