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