diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h
index 95826f9b60..da3b7ba50e 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h
+++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h
@@ -174,6 +174,9 @@ static constexpr int warpSize = 64;
 __device__ long long int clock64();
 __device__ clock_t clock();
 
+//abort
+__device__ void abort();
+
 //atomicAdd()
 __device__ int atomicAdd(int* address, int val);
 __device__ unsigned int atomicAdd(unsigned int* address,
diff --git a/projects/clr/hipamd/src/device_util.cpp b/projects/clr/hipamd/src/device_util.cpp
index 062372f0f4..1efda02933 100644
--- a/projects/clr/hipamd/src/device_util.cpp
+++ b/projects/clr/hipamd/src/device_util.cpp
@@ -839,6 +839,11 @@ __device__ float __hip_ynf(int n, float x)
 __device__ long long int clock64() { return (long long int)hc::__cycle_u64(); };
 __device__ clock_t clock() { return (clock_t)hc::__cycle_u64(); };
 
+//abort
+__device__ void abort()
+{
+    return hc::abort();
+}
 
 //atomicAdd()
 __device__  int atomicAdd(int* address, int val)