Add __assert_fail, __device_trap and hipErrorAssert for clang

include/hip/hcc_detail/hip_runtime.h

@@ -612,6 +612,22 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim;
 #define hipGridDim_y gridDim.y
 #define hipGridDim_z gridDim.z
 
+#pragma push_macro("__DEVICE__")
+#define __DEVICE__ extern "C" __device__ inline __attribute__((always_inline)) \
+  __attribute__((weak))
+
+__DEVICE__ void __device_trap() __asm("llvm.trap");
+
+__DEVICE__ void __assert_fail(const char * __assertion,
+                              const char *__file,
+                              unsigned int __line,
+                              const char *__function)
+{
+    // Ignore all the args for now.
+    __device_trap();
+}
+#pragma push_macro("__DEVICE__")
+
 #endif
 
 #endif  // HIP_HCC_DETAIL_RUNTIME_H

include/hip/hip_runtime_api.h

@@ -243,6 +243,8 @@ typedef enum __HIP_NODISCARD hipError_t {
         1062,  ///< Produced when trying to unlock a non-page-locked memory.
     hipErrorMapBufferObjectFailed =
         1071,    ///< Produced when the IPC memory attach failed from ROCr.
+    hipErrorAssert =
+        1081,    ///< Produced when the kernel calls assert.
     hipErrorTbd  ///< Marker that more error codes are needed.
 } hipError_t;