[compiler-rt] Add SPIR-V target support for UBSan Minimal Runtime #195979
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,76 @@ | ||
|
|
||
| #include "ubsan_minimal_common.h" | ||
|
Contributor
There was a problem hiding this comment. Not a fan of splitting this up so arbitrarily, I don't know what SPIR-V does that can't be done by improving the set of compiler builtins
Contributor
Author
There was a problem hiding this comment. The
Contributor
There was a problem hiding this comment. I guess I don't really understand how the SPIR-V build is supposed to work, but why do we have both C libcalls and stuff like __builtin_spirv_num_workgroups(0)? I generally prefer the latter.
Contributor
Author
There was a problem hiding this comment. I've updated the file to use the |
||
|
|
||
| #include <stdlib.h> | ||
| #include <string.h> | ||
|
|
||
| #if defined(__ANDROID__) | ||
| extern "C" __attribute__((weak)) void android_set_abort_message(const char *); | ||
| #endif // defined(__ANDROID__) | ||
|
|
||
| #ifdef KERNEL_USE | ||
| extern "C" void ubsan_message(const char *msg); | ||
| static void message(const char *msg) { ubsan_message(msg); } | ||
| #elif defined(SANITIZER_AMDGPU) || defined(SANITIZER_NVPTX) | ||
| #include <stdio.h> | ||
| static void message(const char *msg) { fprintf(stderr, "%s", msg); } | ||
| #else | ||
| #include <unistd.h> | ||
| static void message(const char *msg) { (void)write(2, msg, strlen(msg)); } | ||
| #endif | ||
|
|
||
| static char *append_str(const char *s, char *buf, const char *end) { | ||
| for (const char *p = s; (buf < end) && (*p != '\0'); ++p, ++buf) | ||
| *buf = *p; | ||
| return buf; | ||
| } | ||
|
|
||
| static char *append_hex(uintptr_t d, char *buf, const char *end) { | ||
| // Print the address by nibbles. | ||
| for (unsigned shift = sizeof(uintptr_t) * 8; shift && buf < end;) { | ||
| shift -= 4; | ||
| unsigned nibble = (d >> shift) & 0xf; | ||
| *(buf++) = nibble < 10 ? nibble + '0' : nibble - 10 + 'a'; | ||
| } | ||
| return buf; | ||
| } | ||
|
|
||
| static void format_msg(const char *kind, uintptr_t caller, char *buf, | ||
| const char *end) { | ||
| buf = append_str("ubsan: ", buf, end); | ||
| buf = append_str(kind, buf, end); | ||
| buf = append_str(" by 0x", buf, end); | ||
| buf = append_hex(caller, buf, end); | ||
| buf = append_str("\n", buf, end); | ||
| if (buf == end) | ||
| --buf; // Make sure we don't cause a buffer overflow. | ||
| *buf = '\0'; | ||
| } | ||
|
|
||
| void __ubsan_message(const char *msg) { | ||
| message(msg); | ||
| } | ||
|
|
||
| void __ubsan_message(const char *kind, uintptr_t caller) { | ||
| char buf[128]; | ||
| format_msg(kind, caller, buf, buf + sizeof(buf)); | ||
| message(buf); | ||
| } | ||
|
|
||
| void __ubsan_abort() { | ||
| abort(); | ||
| } | ||
|
|
||
| void __ubsan_abort_with_message(const char *kind, uintptr_t caller) { | ||
| char buf[128]; | ||
| format_msg(kind, caller, buf, buf + sizeof(buf)); | ||
|
|
||
| #if defined(__ANDROID__) | ||
| if (&android_set_abort_message) | ||
| android_set_abort_message(buf); | ||
| #else // defined(__ANDROID__) | ||
| message(buf); | ||
| #endif // defined(__ANDROID__) | ||
|
|
||
| abort(); | ||
| } | ||
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,45 @@ | ||
| #include "ubsan_minimal_common.h" | ||
|
|
||
| using uint3 = unsigned __attribute__((ext_vector_type(3))); | ||
|
Contributor
There was a problem hiding this comment. Why is all of this necessary?
Contributor
Author
There was a problem hiding this comment. I believe, that the group ID and local ID info of the error should be useful for the GPU targets. Here is the OpenCL Reference: https://registry.khronos.org/OpenCL/specs/unified/html/OpenCL_C.html#printf-format-string |
||
|
|
||
| struct __ubsan_abort_info_t { | ||
| __attribute__((opencl_constant)) const char *fmt; | ||
| const char *kind; | ||
| uintptr_t caller; | ||
| uint3 gid; | ||
| uint3 lid; | ||
| }; | ||
|
|
||
| // OpenCL printf maps to OpExtInst printf (OpenCL extended instruction set). | ||
| extern "C" int printf(__attribute__((opencl_constant)) const char *fmt, ...); | ||
|
|
||
| // OpenCL work-item builtins map to SPIR-V BuiltIn variables. | ||
| extern "C" unsigned get_group_id(unsigned dim); | ||
| extern "C" unsigned get_local_id(unsigned dim); | ||
|
|
||
| static __attribute__((opencl_constant)) const char ubsan_msg_simple[] = "%s"; | ||
| static __attribute__((opencl_constant)) const char ubsan_msg_fmt[] = | ||
| "ubsan: %s by 0x%lx at gid=[%v3u] lid=[%v3u]\n"; | ||
|
|
||
| void __ubsan_message(const char *msg) { printf(ubsan_msg_simple, msg); } | ||
|
|
||
| void __ubsan_message(const char *kind, uintptr_t caller) { | ||
| uint3 gid = {get_group_id(0), get_group_id(1), get_group_id(2)}; | ||
| uint3 lid = {get_local_id(0), get_local_id(1), get_local_id(2)}; | ||
|
|
||
| printf(ubsan_msg_fmt, kind, caller, gid, lid); | ||
| } | ||
|
|
||
| // SPV_KHR_abort: OpAbortKHR terminates the invocation and passes a message | ||
| // to the client API. The message is passed as a typed value. | ||
| [[noreturn]] void __spirv_AbortKHR(__ubsan_abort_info_t info); | ||
|
Contributor
There was a problem hiding this comment. This should not be defined here, it should be what
Contributor
Author
There was a problem hiding this comment. The See the SPIR-V extension spec for more info: https://github.khronos.org/SPIRV-Registry/extensions/KHR/SPV_KHR_abort.html
Contributor
There was a problem hiding this comment. Trapping isn't the intended way for the minimal runtime to work. If you just want to abort on first error you'd use
Contributor
Author
There was a problem hiding this comment. The idea is to report an error and continue, when the error is recoverable. There are two kinds of the error handlers: non-aborting and aborting. For the non-aborting handlers we just should the Also, the |
||
|
|
||
| [[noreturn]] void __ubsan_abort() { __ubsan_abort_with_message("abort", 0); } | ||
|
|
||
| [[noreturn]] void __ubsan_abort_with_message(const char *kind, | ||
| uintptr_t caller) { | ||
| uint3 gid = {get_group_id(0), get_group_id(1), get_group_id(2)}; | ||
| uint3 lid = {get_local_id(0), get_local_id(1), get_local_id(2)}; | ||
| __ubsan_abort_info_t info = {ubsan_msg_fmt, kind, caller, gid, lid}; | ||
| __spirv_AbortKHR(info); | ||
| } | ||
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,32 @@ | ||
| //===-- ubsan_minimal_common.h ----------------------------------*- C++ -*-===// | ||
| // | ||
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | ||
| // See https://llvm.org/LICENSE.txt for license information. | ||
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
| // | ||
| // This file is a part of the minimal UB sanitizer runtime. | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
|
|
||
| #ifndef UBSAN_MINIMAL_COMMON_H | ||
| #define UBSAN_MINIMAL_COMMON_H | ||
|
|
||
| #if defined(__UINTPTR_TYPE__) | ||
| typedef __UINTPTR_TYPE__ uintptr_t; | ||
|
|
||
| static_assert(sizeof(uintptr_t) == sizeof(void *), | ||
| "uintptr_t must be the same size as void*"); | ||
| #else // defined(__UINTPTR_TYPE__) | ||
| #include <stdint.h> | ||
| #endif // defined(__UINTPTR_TYPE__) | ||
|
|
||
| void __ubsan_message(const char *msg); | ||
| void __ubsan_message(const char *kind, uintptr_t caller); | ||
|
|
||
| [[noreturn]] void __ubsan_abort(); | ||
| [[noreturn]] void __ubsan_abort_with_message(const char *kind, | ||
| uintptr_t caller); | ||
|
|
||
| #endif // UBSAN_MINIMAL_COMMON_H |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Isn't this case handled above?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Fixed: 06ffdd6