AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-234709 - SWDEV-258907 - Reimplement assert() using the new fprintf_stderr facility

Browse Source 
This reverts commit 18234e1d24.

Change-Id: Ic1bbf4e4e6f950934d9ba02224ef61a96bdd0e9b
This commit is contained in:
Maneesh Gupta
2021-07-07 00:55:38 -07:00
parent 67be27c184
commit 1e35b26108
2 changed files with 42 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/include/hip/amd_detail/amd_device_functions.h
+34 -3
View File
@@ -1175,9 +1175,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
+8
View File
@@ -81,6 +81,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)))