HIP: Heterogenous-computing Interface for Portability
hip_assert.h
1 /*
2 Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
3 
4 Permission is hereby granted, free of charge, to any person obtaining a copy
5 of this software and associated documentation files (the "Software"), to deal
6 in the Software without restriction, including without limitation the rights
7 to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
8 copies of the Software, and to permit persons to whom the Software is
9 furnished to do so, subject to the following conditions:
10 
11 The above copyright notice and this permission notice shall be included in
12 all copies or substantial portions of the Software.
13 
14 THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15 IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16 FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17 AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18 LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
19 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
20 THE SOFTWARE.
21 */
22 
23 #pragma once
24 
25 // abort
26 extern "C" __device__ inline __attribute__((weak))
27 void 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)
37 extern "C" __device__ __attribute__((noinline)) __attribute__((weak))
38 void _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) */
43 extern "C" __device__ __attribute__((noinline)) __attribute__((weak))
44 void __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 
85 extern "C" __device__ __attribute__((noinline)) __attribute__((weak))
86 void __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
_Float16 __2f16 __attribute__((ext_vector_type(2)))
Definition: hip_fp16_math_fwd.h:57