Author: Artem Belevich Date: 2021-03-01T13:59:22-08:00 New Revision: 32e0645276230bb5b736e378860df3b92b1f4ba8
URL: https://github.com/llvm/llvm-project/commit/32e0645276230bb5b736e378860df3b92b1f4ba8 DIFF: https://github.com/llvm/llvm-project/commit/32e0645276230bb5b736e378860df3b92b1f4ba8.diff LOG: [CUDA] Remove `noreturn` attribute from __assertfail(). `noreturn` complicates control flow and tends to trigger a known bug in ptxas if the assert is used within loops in sufficiently complicated code. https://bugs.llvm.org/show_bug.cgi?id=27738 Differential Revision: https://reviews.llvm.org/D97708 Added: Modified: clang/lib/Headers/__clang_cuda_runtime_wrapper.h Removed: ################################################################################ diff --git a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h index f88c39a9b6e5..f401964bd529 100644 --- a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h +++ b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h @@ -349,9 +349,14 @@ extern "C" { __device__ int vprintf(const char *, const char *); __device__ void free(void *) __attribute((nothrow)); __device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc)); + +// __assertfail() used to have a `noreturn` attribute. Unfortunately that +// contributed to triggering the longstanding bug in ptxas when assert was used +// in sufficiently convoluted code. See +// https://bugs.llvm.org/show_bug.cgi?id=27738 for the details. __device__ void __assertfail(const char *__message, const char *__file, unsigned __line, const char *__function, - size_t __charSize) __attribute__((noreturn)); + size_t __charSize); // In order for standard assert() macro on linux to work we need to // provide device-side __assert_fail() _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits