[clang] [CUDA][HIP] Add a __device__ version of std::__glibcxx_assert_fail() (PR #136133)

[Juan Manuel Martinez Caamaño via cfe-commits]

jmmartinez wrote:

I've added a verbose version of the assertion that uses `__builtin_printf` (I could not figure out a way to correctly emit something to `stderr` from device code in a generic way without including `cstdio`.

```cpp
#ifdef _GLIBCXX_VERBOSE_ASSERT
__attribute__((device, noreturn)) inline void
__glibcxx_assert_fail(const char *file, int line, const char *function,
                      const char *condition) noexcept {
  if (file && function && condition)
    __builtin_printf("%s:%d: %s: Assertion '%s' failed.\n", file, line,
                     function, condition);
  else if (function)
    __builtin_printf("%s: Undefined behavior detected.\n", function);
  __builtin_abort();
}
#endif
```

An example from out test with an out-of-bounds access to an `std::array` (the first line is from the assert, the rest comes from the exception raised):

```cpp
...
/home/juamarti/swdev/518041/_gcc/_install/lib/gcc/x86_64-pc-linux-gnu/15.0.1/../../../../include/c++/15.0.1/array:210: reference std::array<float, 16>::operator[](size_type) [_Tp = float, _Nm = 16]: Assertion '__n < this->size()' failed.
Kernel Name: _Z15vectoradd_floatPfPKfS1_
VGPU=0x5bae1d0e41e0 SWq=0x7929b05da000, HWq=0x792897b00000, id=1                                                                                                                                                         Dispatch Header = 0xb02 (type=2, barrier=1, acquire=1, release=1), setup=0
        grid=[65536, 1, 1], workgroup=[16, 1, 1]                                                                                                                                                                         private_seg_size=1920, group_seg_size=0
        kernel_obj=0x7929aea90ac0, kernarg_address=0x0x792895d00000                                                                                                                                                      completion_signal=0x0, correlation_id=0
        rptr=4, wptr=6
:0:rocdevice.cpp            :3620: 8952765750346 us:  Callback: Queue 0x792897b00000 aborting with error : HSA_STATUS_ERROR_EXCEPTION: An HSAIL operation resulted in a hardware exception. code: 0x1016
Aborted (core dumped)
```

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