tra updated this revision to Diff 327197. tra edited the summary of this revision. tra added a comment.
Added a comment. Repository: rG LLVM Github Monorepo CHANGES SINCE LAST ACTION https://reviews.llvm.org/D97708/new/ https://reviews.llvm.org/D97708 Files: clang/lib/Headers/__clang_cuda_runtime_wrapper.h Index: clang/lib/Headers/__clang_cuda_runtime_wrapper.h =================================================================== --- clang/lib/Headers/__clang_cuda_runtime_wrapper.h +++ clang/lib/Headers/__clang_cuda_runtime_wrapper.h @@ -349,9 +349,14 @@ __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()
Index: clang/lib/Headers/__clang_cuda_runtime_wrapper.h =================================================================== --- clang/lib/Headers/__clang_cuda_runtime_wrapper.h +++ clang/lib/Headers/__clang_cuda_runtime_wrapper.h @@ -349,9 +349,14 @@ __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
