From 2cac768a07aeea2dc35244e378afa97bc711278c Mon Sep 17 00:00:00 2001 From: Sameer Sahasrabuddhe Date: Fri, 15 Jan 2021 11:20:40 +0530 Subject: [PATCH] Implement assert() using the new fprintf_stderr facility The device library now provides functions to specify the host stderr stream in a hostcall printf message. This change implements an assert() macro which can construct such a message. The end result is that assertions on the device are now correctly printed on the host stderr instead of stdout. Change-Id: I85ab8f7848bcf28303cb8dbb8a798bc6aece7d75 --- .../include/hip/amd_detail/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/device_functions.h b/hipamd/include/hip/amd_detail/device_functions.h index fa8844d331..320fbc7422 100644 --- a/hipamd/include/hip/amd_detail/device_functions.h +++ b/hipamd/include/hip/amd_detail/device_functions.h @@ -1117,9 +1117,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 a8fb2deecc..7e021bfa68 100644 --- a/hipamd/include/hip/amd_detail/device_library_decls.h +++ b/hipamd/include/hip/amd_detail/device_library_decls.h @@ -78,6 +78,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)))