[clang] 32e0645 - [CUDA] Remove `noreturn` attribute from __assertfail().

[Artem Belevich via cfe-commits]

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()