From 40ff42e555e5c6c3eb7f54b7db8237f267b3fced Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Thu, 7 Sep 2017 20:45:27 +0530
Subject: [PATCH] Device side assert support

---
 include/hip/hcc_detail/hip_runtime.h  | 2 +-
 include/hip/nvcc_detail/hip_runtime.h | 2 ++
 2 files changed, 3 insertions(+), 1 deletion(-)

diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h
index 1092e22c47..2386ea08cb 100644
--- a/include/hip/hcc_detail/hip_runtime.h
+++ b/include/hip/hcc_detail/hip_runtime.h
@@ -99,7 +99,7 @@ extern int HIP_TRACE_API;
 // TODO-HCC add a dummy implementation of assert, need to replace with a proper kernel exit call.
 #if __HIP_DEVICE_COMPILE__ == 1
    #undef assert
-   #define assert(COND) { if (COND) {} }
+   #define assert(COND) { if (!COND) {abort();} }
 #endif
 
 
diff --git a/include/hip/nvcc_detail/hip_runtime.h b/include/hip/nvcc_detail/hip_runtime.h
index 8c08f3d151..fca9ab6e39 100644
--- a/include/hip/nvcc_detail/hip_runtime.h
+++ b/include/hip/nvcc_detail/hip_runtime.h
@@ -111,6 +111,8 @@ kernelName<<<numblocks,numthreads,memperblock,streamId>>>(__VA_ARGS__);\
 
 #ifdef __HIP_DEVICE_COMPILE__
 #define abort() {asm("trap;");}
+#undef assert
+#define assert(COND) { if (!COND) {abort();} }
 #endif
 
 #endif