From 7672b44c790eed7cefdf8d93a2caaeb86a7ac190 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu@amd.com>
Date: Tue, 8 May 2018 15:36:01 -0400
Subject: [PATCH] Add __assert_fail, __device_trap and hipErrorAssert for clang

---
 include/hip/hcc_detail/hip_runtime.h | 16 ++++++++++++++++
 include/hip/hip_runtime_api.h        |  2 ++
 2 files changed, 18 insertions(+)

diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h
index 92f06e9174..e9551d4f9a 100644
--- a/include/hip/hcc_detail/hip_runtime.h
+++ b/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
diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h
index bfde2e942d..2b36f3e140 100644
--- a/include/hip/hip_runtime_api.h
+++ b/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;