SWDEV-269351 - Revert "Implement assert() using the new fprintf_stderr facility"

This reverts commit 96be8a7b73.

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

@@ -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))

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)))