HIP: Heterogenous-computing Interface for Portability
Loading...
Searching...
No Matches
hip_assert.h
1/*
2Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
3
4Permission is hereby granted, free of charge, to any person obtaining a copy
5of this software and associated documentation files (the "Software"), to deal
6in the Software without restriction, including without limitation the rights
7to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
8copies of the Software, and to permit persons to whom the Software is
9furnished to do so, subject to the following conditions:
10
11The above copyright notice and this permission notice shall be included in
12all copies or substantial portions of the Software.
13
14THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
19OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
20THE SOFTWARE.
21*/
22
23#pragma once
24
25// abort
26extern "C" __device__ inline __attribute__((weak))
27void abort() {
28 __builtin_trap();
29}
30
31// The noinline attribute helps encapsulate the printf expansion,
32// which otherwise has a performance impact just by increasing the
33// size of the calling function. Additionally, the weak attribute
34// allows the function to exist as a global although its definition is
35// included in every compilation unit.
36#if defined(_WIN32) || defined(_WIN64)
37extern "C" __device__ __attribute__((noinline)) __attribute__((weak))
38void _wassert(const wchar_t *_msg, const wchar_t *_file, unsigned _line) {
39 // FIXME: Need `wchar_t` support to generate assertion message.
40 __builtin_trap();
41}
42#else /* defined(_WIN32) || defined(_WIN64) */
43extern "C" __device__ __attribute__((noinline)) __attribute__((weak))
44void __assert_fail(const char *assertion,
45 const char *file,
46 unsigned int line,
47 const char *function)
48{
49 const char fmt[] = "%s:%u: %s: Device-side assertion `%s' failed.\n";
50
51 // strlen is not available as a built-in yet, so we create our own
52 // loop in a macro. With a string literal argument, the compiler
53 // usually manages to replace the loop with a constant.
54 //
55 // The macro does not check for null pointer, since all the string
56 // arguments are defined to be constant literals when called from
57 // the assert() macro.
58 //
59 // NOTE: The loop below includes the null terminator in the length
60 // as required by append_string_n().
61#define __hip_get_string_length(LEN, STR) \
62 do { \
63 const char *tmp = STR; \
64 while (*tmp++); \
65 LEN = tmp - STR; \
66 } while (0)
67
68 auto msg = __ockl_fprintf_stderr_begin();
69 int len = 0;
70 __hip_get_string_length(len, fmt);
71 msg = __ockl_fprintf_append_string_n(msg, fmt, len, 0);
72 __hip_get_string_length(len, file);
73 msg = __ockl_fprintf_append_string_n(msg, file, len, 0);
74 msg = __ockl_fprintf_append_args(msg, 1, line, 0, 0, 0, 0, 0, 0, 0);
75 __hip_get_string_length(len, function);
76 msg = __ockl_fprintf_append_string_n(msg, function, len, 0);
77 __hip_get_string_length(len, assertion);
78 __ockl_fprintf_append_string_n(msg, assertion, len, /* is_last = */ 1);
79
80#undef __hip_get_string_length
81
82 __builtin_trap();
83}
84
85extern "C" __device__ __attribute__((noinline)) __attribute__((weak))
86void __assertfail()
87{
88 // ignore all the args for now.
89 __builtin_trap();
90}
91#endif /* defined(_WIN32) || defined(_WIN64) */
92
93#if defined(NDEBUG)
94#define __hip_assert(COND)
95#else
96#define __hip_assert(COND) \
97 do { \
98 if (!(COND)) \
99 __builtin_trap(); \
100 } while (0)
101#endif