From 7998429c2102f8338aec4fa16a2dd4437666b556 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 12 Feb 2021 03:44:14 -0500 Subject: [PATCH] SWDEV-269351 - Revert "Implement assert() using the new fprintf_stderr facility" This reverts commit 96be8a7b73c4d5d21085e5c68211c8117a55e1ba. Reverting since this is dependent on a device libs change which is not yet in mainline. Change-Id: I094781b868f39995d1eecb41488c1a1a0147b787 --- include/hip/amd_detail/device_functions.h | 37 ++----------------- include/hip/amd_detail/device_library_decls.h | 8 ---- 2 files changed, 3 insertions(+), 42 deletions(-) diff --git a/include/hip/amd_detail/device_functions.h b/include/hip/amd_detail/device_functions.h index 273912256b..e476a1a013 100644 --- a/include/hip/amd_detail/device_functions.h +++ b/include/hip/amd_detail/device_functions.h @@ -1117,40 +1117,9 @@ void __assert_fail(const char *assertion, unsigned int line, const char *function) { - 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(); + printf("%s:%u: %s: Device-side assertion `%s' failed.\n", file, line, + function, assertion); + __builtin_trap(); } extern "C" __device__ __attribute__((noinline)) __attribute__((weak)) diff --git a/include/hip/amd_detail/device_library_decls.h b/include/hip/amd_detail/device_library_decls.h index 7e021bfa68..a8fb2deecc 100644 --- a/include/hip/amd_detail/device_library_decls.h +++ b/include/hip/amd_detail/device_library_decls.h @@ -78,14 +78,6 @@ 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)))