diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h
index 4cbaf9ea4e..135e307b6e 100644
--- a/include/hip/hcc_detail/hip_runtime_api.h
+++ b/include/hip/hcc_detail/hip_runtime_api.h
@@ -60,6 +60,7 @@ typedef void* hipDeviceptr_t;
 
 typedef struct ihipEvent_t *hipEvent_t;
 
+typedef unsigned hipLimit;
 
 /**
  * @addtogroup GlobalDefs More
@@ -97,6 +98,8 @@ typedef struct ihipEvent_t *hipEvent_t;
 #define hipDeviceMapHost            0x8
 #define hipDeviceLmemResizeToMax    0x16
 
+#define hipLimitMallocHeapSize      0x2
+
 /**
  * @warning On AMD devices and recent Nvidia devices, these hints and controls are ignored.
  */
@@ -323,6 +326,18 @@ hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig );
  */
 hipError_t hipDeviceGetCacheConfig ( hipFuncCache *cacheConfig );
 
+/**
+ * @brief Get Resource limits of current device
+ *
+ * @param [out] pValue
+ * @param [in]  limit
+ *
+ * @returns #hipSuccess, #hipErrorUnsupportedLimit, #hipErrorInvalidValue
+ * Note: Currently, only hipLimitMallocHeapSize is available
+ *
+ */
+hipError_t hipDeviceGetLimit(size_t *pValue, hipLimit limit);
+
 
 /**
  * @brief Set Cache configuration for a specific function
diff --git a/src/hip_device.cpp b/src/hip_device.cpp
index a677402b69..db04985cf1 100644
--- a/src/hip_device.cpp
+++ b/src/hip_device.cpp
@@ -68,7 +68,7 @@ hipError_t hipGetDeviceCount(int *count)
     return e;
 }
 
-hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig )
+hipError_t hipDeviceSetCacheConfig(hipFuncCache cacheConfig)
 {
     HIP_INIT_API(cacheConfig);
 
@@ -77,7 +77,7 @@ hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig )
     return ihipLogStatus(hipSuccess);
 }
 
-hipError_t hipDeviceGetCacheConfig ( hipFuncCache *cacheConfig )
+hipError_t hipDeviceGetCacheConfig(hipFuncCache *cacheConfig)
 {
     HIP_INIT_API(cacheConfig);
 
@@ -86,7 +86,23 @@ hipError_t hipDeviceGetCacheConfig ( hipFuncCache *cacheConfig )
     return ihipLogStatus(hipSuccess);
 }
 
-hipError_t hipFuncSetCacheConfig ( hipFuncCache cacheConfig )
+extern "C" size_t g_malloc_heap_size;
+
+hipError_t hipDeviceGetLimit (size_t *pValue, hipLimit limit)
+{
+    HIP_INIT_API(pValue, limit);
+    if(pValue == nullptr) {
+        return ihipLogStatus(hipErrorInvalidValue);
+    }
+    if(limit == hipLimitMallocHeapSize) {
+        *pValue = g_malloc_heap_size;
+        return ihipLogStatus(hipSuccess);
+    }else{
+        return ihipLogStatus(hipErrorUnsupportedLimit);
+    }
+}
+
+hipError_t hipFuncSetCacheConfig (hipFuncCache cacheConfig)
 {
     HIP_INIT_API(cacheConfig);
 
@@ -95,7 +111,7 @@ hipError_t hipFuncSetCacheConfig ( hipFuncCache cacheConfig )
     return ihipLogStatus(hipSuccess);
 }
 
-hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config )
+hipError_t hipDeviceSetSharedMemConfig (hipSharedMemConfig config)
 {
     HIP_INIT_API(config);
 
@@ -104,7 +120,7 @@ hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config )
     return ihipLogStatus(hipSuccess);
 }
 
-hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig )
+hipError_t hipDeviceGetSharedMemConfig (hipSharedMemConfig *pConfig)
 {
     HIP_INIT_API(pConfig);
 
diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp
index 06f06a65b3..0ba98b8a4c 100644
--- a/src/hip_hcc.cpp
+++ b/src/hip_hcc.cpp
@@ -86,8 +86,6 @@ hsa_agent_t g_cpu_agent;
 
  This is the best place to put them because the device
  global variables need to be initialized at the start.
-
-
 */
 
 #define NUM_PAGES_PER_THREAD  16
@@ -98,47 +96,44 @@ hsa_agent_t g_cpu_agent;
 #define SIZE_MALLOC NUM_PAGES * SIZE_OF_PAGE
 #define SIZE_OF_HEAP SIZE_MALLOC
 
-struct heapTracker_t {
-  void *ptr;
-  uint32_t *flags;
-  uint32_t next;
-};
+size_t g_malloc_heap_size = SIZE_OF_HEAP;
 
 __attribute__((address_space(1))) char gpuHeap[SIZE_OF_HEAP];
 __attribute__((address_space(1))) uint32_t gpuFlags[NUM_PAGES];
 
-__device__ void *__hip_hc_malloc(size_t size){
-  char *heap = (char*)gpuHeap;
-  if(size > SIZE_OF_HEAP)
-  {
-      return (void*)nullptr;
-  }
-  uint32_t totalThreads = hipBlockDim_x * hipGridDim_x * hipBlockDim_y * hipGridDim_y * hipBlockDim_z * hipGridDim_z;
-  uint32_t currentWorkItem = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x;
+__device__ void *__hip_hc_malloc(size_t size) 
+{
+    char *heap = (char*)gpuHeap;
+    if(size > SIZE_OF_HEAP)
+    {
+        return (void*)nullptr;
+    }
+    uint32_t totalThreads = hipBlockDim_x * hipGridDim_x * hipBlockDim_y * hipGridDim_y * hipBlockDim_z * hipGridDim_z;
+    uint32_t currentWorkItem = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x;
 
-  uint32_t numHeapsPerWorkItem = NUM_PAGES / totalThreads;
-  uint32_t heapSizePerWorkItem = SIZE_OF_HEAP / totalThreads;
+    uint32_t numHeapsPerWorkItem = NUM_PAGES / totalThreads;
+    uint32_t heapSizePerWorkItem = SIZE_OF_HEAP / totalThreads;
 
-  uint32_t stride = size / SIZE_OF_PAGE;
-  uint32_t start = numHeapsPerWorkItem * currentWorkItem;
+    uint32_t stride = size / SIZE_OF_PAGE;
+    uint32_t start = numHeapsPerWorkItem * currentWorkItem;
 
-  uint32_t k=0;
+    uint32_t k=0;
 
-  while(gpuFlags[k] > 0)
-  {
-      k++;
-  }
+    while(gpuFlags[k] > 0)
+    {
+        k++;
+    }
 
-  for(uint32_t i=0;i<stride-1;i++)
-  {
-      gpuFlags[i+start+k] = 1;
-  }
+    for(uint32_t i=0;i<stride-1;i++)
+    {
+        gpuFlags[i+start+k] = 1;
+    }
 
-  gpuFlags[start+stride-1+k] = 2;
+    gpuFlags[start+stride-1+k] = 2;
 
-  void *ptr = (void*)(heap + heapSizePerWorkItem * currentWorkItem + k*SIZE_OF_PAGE);
+    void* ptr = (void*)(heap + heapSizePerWorkItem * currentWorkItem + k*SIZE_OF_PAGE);
 
-  return ptr;
+    return ptr;
 }
 
 __device__ void* __hip_hc_free(void *ptr)
@@ -151,15 +146,14 @@ __device__ void* __hip_hc_free(void *ptr)
     uint32_t offsetByte = (uint64_t)ptr - (uint64_t)gpuHeap;
     uint32_t offsetPage = offsetByte / SIZE_OF_PAGE;
 
-    while(gpuFlags[offsetPage] != 0)
-    {
-        if(gpuFlags[offsetPage] == 2){
-          gpuFlags[offsetPage] = 0;
-          offsetPage++;
-          break;
-        }else{
-          gpuFlags[offsetPage] = 0;
-          offsetPage++;
+    while(gpuFlags[offsetPage] != 0) {
+        if(gpuFlags[offsetPage] == 2) {
+            gpuFlags[offsetPage] = 0;
+            offsetPage++;
+            break;
+        } else {
+            gpuFlags[offsetPage] = 0;
+            offsetPage++;
         }
     }
 
diff --git a/tests/src/deviceLib/hipTestDeviceLimit.cpp b/tests/src/deviceLib/hipTestDeviceLimit.cpp
new file mode 100644
index 0000000000..076bfdc61b
--- /dev/null
+++ b/tests/src/deviceLib/hipTestDeviceLimit.cpp
@@ -0,0 +1,10 @@
+#include<hip/hip_runtime_api.h>
+#include<iostream>
+#include<assert.h>
+
+int main()
+{
+    size_t heap;
+    assert(hipSuccess == hipDeviceGetLimit(&heap, hipLimitMallocHeapSize));
+    assert(heap == 4194304);
+}