@@ -359,6 +359,7 @@ static inline int C10_WARP_SIZE_INTERNAL() {
359359// Those platforms do not support assert()
360360#define CUDA_KERNEL_ASSERT (cond )
361361#define CUDA_KERNEL_ASSERT_MSG (cond, msg )
362+ #define CUDA_KERNEL_ASSERT_PRINTF (cond, msg, ...)
362363#define SYCL_KERNEL_ASSERT (cond )
363364#elif defined(_MSC_VER)
364365#if defined(NDEBUG)
@@ -396,6 +397,26 @@ __host__ __device__
396397 static_cast <unsigned >(__LINE__)), \
397398 0 ); \
398399 }
400+ #define CUDA_KERNEL_ASSERT_PRINTF (cond, msg, ...) \
401+ if (C10_UNLIKELY (!(cond))) { \
402+ (void )(printf ( \
403+ " [CUDA_KERNEL_ASSERT] " __FILE__ " :" C10_STRINGIZE ( \
404+ __LINE__) " : %s: block: [%d,%d,%d], thread: [%d,%d,%d]: " \
405+ " Assertion failed: `" #cond " `: " msg " \n " , \
406+ __func__, \
407+ blockIdx.x , \
408+ blockIdx.y , \
409+ blockIdx.z , \
410+ threadIdx.x , \
411+ threadIdx.y , \
412+ threadIdx.z , \
413+ ##__VA_ARGS__)); \
414+ (void )(_wassert ( \
415+ _CRT_WIDE (#cond), \
416+ _CRT_WIDE (__FILE__), \
417+ static_cast <unsigned >(__LINE__)), \
418+ 0 ); \
419+ }
399420#define SYCL_KERNEL_ASSERT (cond ) \
400421 if (C10_UNLIKELY (!(cond))) { \
401422 (void )(_wassert ( \
@@ -455,6 +476,10 @@ __host__ __device__
455476 if C10_UNLIKELY (!(cond)) { \
456477 abort (); \
457478 }
479+ #define CUDA_KERNEL_ASSERT_PRINTF (cond, msg, ...) \
480+ if C10_UNLIKELY (!(cond)) { \
481+ abort (); \
482+ }
458483#define SYCL_KERNEL_ASSERT (cond ) \
459484 if C10_UNLIKELY (!(cond)) { \
460485 abort (); \
@@ -470,6 +495,23 @@ __host__ __device__
470495 __assert_fail ( \
471496 msg, __FILE__, static_cast <unsigned int >(__LINE__), __func__); \
472497 }
498+ #define CUDA_KERNEL_ASSERT_PRINTF (cond, msg, ...) \
499+ if (C10_UNLIKELY (!(cond))) { \
500+ printf ( \
501+ " [CUDA_KERNEL_ASSERT] " __FILE__ " :" C10_STRINGIZE ( \
502+ __LINE__) " : %s: block: [%d,%d,%d], thread: [%d,%d,%d]: " \
503+ " Assertion failed: `" #cond " `: " msg " \n " , \
504+ __func__, \
505+ blockIdx.x , \
506+ blockIdx.y , \
507+ blockIdx.z , \
508+ threadIdx.x , \
509+ threadIdx.y , \
510+ threadIdx.z , \
511+ ##__VA_ARGS__); \
512+ __assert_fail ( \
513+ #cond, __FILE__, static_cast <unsigned int >(__LINE__), __func__); \
514+ }
473515#define SYCL_KERNEL_ASSERT (cond ) \
474516 if (C10_UNLIKELY (!(cond))) { \
475517 __assert_fail ( \
0 commit comments