[CUDA] Remove `noreturn` attribute from __assertfail().

[tra]

Summary

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

[jlebar]

Maybe add a comment on assertfail that we don't want it to be noreturn because blah blah? So someone doesn't come back and add it when we're not looking. :)

[tra]

Added a comment.

clang/lib/Headers/__clang_cuda_runtime_wrapper.h

 // Device-side CUDA system calls.
 // http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls
 // We need these declarations and wrappers for device-side
 // malloc/free/printf calls to work without relying on
 // -fcuda-disable-target-call-checks option.
 __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()
 __device__ static inline void __assert_fail(const char *__message,
                                             const char *__file, unsigned __line,
                                             const char *__function) {
   __assertfail(__message, __file, __line, __function, sizeof(char));
 }