[compiler-rt] [compiler-rt] Add SPIR-V target support for UBSan Minimal Runtime (PR #195979)

[Victor Mustya via llvm-commits]

================
@@ -0,0 +1,45 @@
+#include "ubsan_minimal_common.h"
+
+using uint3 = unsigned __attribute__((ext_vector_type(3)));
+
+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);
----------------
vmustya wrote:

The `__builtin_verbose_trap` doesn't produce any semantic info in the SPIR-V (or any other generated code). It generates the `llvm.trap` intrinsic and some extra debug info. However, the `__spirv_AbortKHR` is lowered into the `OpAbortKHR` SPIR-V instruction with an arbitrary error information. Here the sanitizer message is passed as the abort operand, so the offload device compiler would be able to print it before the kernel termination.

https://github.com/llvm/llvm-project/pull/195979