From 1e35b26108ef3edab2b00596bc5927b0e9cfb589 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 7 Jul 2021 00:55:38 -0700 Subject: [PATCH] SWDEV-234709 - SWDEV-258907 - Reimplement assert() using the new fprintf_stderr facility This reverts commit 18234e1d24f3807a49f5a1c976dc1c990afdfc1b. Change-Id: Ic1bbf4e4e6f950934d9ba02224ef61a96bdd0e9b --- .../hip/amd_detail/amd_device_functions.h | 37 +++++++++++++++++-- .../hip/amd_detail/device_library_decls.h | 8 ++++ 2 files changed, 42 insertions(+), 3 deletions(-) diff --git a/hipamd/include/hip/amd_detail/amd_device_functions.h b/hipamd/include/hip/amd_detail/amd_device_functions.h index 856a3bf829..f51d86c199 100644 --- a/hipamd/include/hip/amd_detail/amd_device_functions.h +++ b/hipamd/include/hip/amd_detail/amd_device_functions.h @@ -1175,9 +1175,40 @@ void __assert_fail(const char *assertion, unsigned int line, const char *function) { - printf("%s:%u: %s: Device-side assertion `%s' failed.\n", file, line, - function, assertion); - __builtin_trap(); + const char fmt[] = "%s:%u: %s: Device-side assertion `%s' failed.\n"; + + // strlen is not available as a built-in yet, so we create our own + // loop in a macro. With a string literal argument, the compiler + // usually manages to replace the loop with a constant. + // + // The macro does not check for null pointer, since all the string + // arguments are defined to be constant literals when called from + // the assert() macro. + // + // NOTE: The loop below includes the null terminator in the length + // as required by append_string_n(). +#define __hip_get_string_length(LEN, STR) \ + do { \ + const char *tmp = STR; \ + while (*tmp++); \ + LEN = tmp - STR; \ + } while (0) + + auto msg = __ockl_fprintf_stderr_begin(); + int len = 0; + __hip_get_string_length(len, fmt); + msg = __ockl_fprintf_append_string_n(msg, fmt, len, 0); + __hip_get_string_length(len, file); + msg = __ockl_fprintf_append_string_n(msg, file, len, 0); + msg = __ockl_fprintf_append_args(msg, 1, line, 0, 0, 0, 0, 0, 0, 0); + __hip_get_string_length(len, function); + msg = __ockl_fprintf_append_string_n(msg, function, len, 0); + __hip_get_string_length(len, assertion); + __ockl_fprintf_append_string_n(msg, assertion, len, /* is_last = */ 1); + +#undef __hip_get_string_length + + __builtin_trap(); } extern "C" __device__ __attribute__((noinline)) __attribute__((weak)) diff --git a/hipamd/include/hip/amd_detail/device_library_decls.h b/hipamd/include/hip/amd_detail/device_library_decls.h index bfb922ff8c..29c9ef4cba 100644 --- a/hipamd/include/hip/amd_detail/device_library_decls.h +++ b/hipamd/include/hip/amd_detail/device_library_decls.h @@ -81,6 +81,14 @@ extern "C" __device__ __attribute__((convergent)) int __ockl_wgred_add_i32(int a extern "C" __device__ __attribute__((convergent)) int __ockl_wgred_and_i32(int a); extern "C" __device__ __attribute__((convergent)) int __ockl_wgred_or_i32(int a); +extern "C" __device__ uint64_t __ockl_fprintf_stderr_begin(); +extern "C" __device__ uint64_t __ockl_fprintf_append_args(uint64_t msg_desc, uint32_t num_args, + uint64_t value0, uint64_t value1, + uint64_t value2, uint64_t value3, + uint64_t value4, uint64_t value5, + uint64_t value6, uint32_t is_last); +extern "C" __device__ uint64_t __ockl_fprintf_append_string_n(uint64_t msg_desc, const char* data, + uint64_t length, uint32_t is_last); // Introduce local address space #define __local __attribute__((address_space(3)))