frozenleaves/pytorch
1
0
Fork 0
mirror of https://github.com/pytorch/pytorch.git synced 2025-10-20 21:14:14 +08:00
Code Issues Packages Projects Releases Wiki Activity

[ROCm] support CUDA_KERNEL_ASSERT using abort() (#155262)

Browse Source 
We won't have the full message that __assert_fail would provide, but at least we won't silently do nothing.

Fixes #155045.

Pull Request resolved: https://github.com/pytorch/pytorch/pull/155262
Approved by: https://github.com/hongxiayang, https://github.com/jeffdaily

Co-authored-by: Jeff Daily <jeff.daily@amd.com>
This commit is contained in:
Jeff Daily
2025-06-18 23:52:35 +00:00
committed by PyTorch MergeBot
parent b8c2d4c259
commit 6303cc41b7
1 changed files with 17 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

21
c10/macros/Macros.h
View File

@ -403,11 +403,24 @@ __host__ __device__
#endif // __SYCL_DEVICE_ONLY__
}
#endif // NDEBUG
// ROCm disable kernel assert by default
// ROCm disables kernel assert by default for performance considerations.
// Though ROCm supports __assert_fail, it uses kernel printf which has
// a non-negligible performance impact even if the assert condition is
// never triggered. We choose to use abort() instead which will still
// terminate the application but without a more useful error message.
#if !defined(C10_USE_ROCM_KERNEL_ASSERT) and defined(USE_ROCM)
#define CUDA_KERNEL_ASSERT(cond)
#define CUDA_KERNEL_ASSERT_MSG(cond, msg)
#define SYCL_KERNEL_ASSERT(cond)
#define CUDA_KERNEL_ASSERT(cond) \
if C10_UNLIKELY (!(cond)) { \
abort(); \
}
#define CUDA_KERNEL_ASSERT_MSG(cond, msg) \
if C10_UNLIKELY (!(cond)) { \
abort(); \
}
#define SYCL_KERNEL_ASSERT(cond) \
if C10_UNLIKELY (!(cond)) { \
abort(); \
}
#else
#define CUDA_KERNEL_ASSERT(cond) \
if (C10_UNLIKELY(!(cond))) { \