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

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

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