diff --git a/include/hip/hcc_detail/hip_memory.h b/include/hip/hcc_detail/hip_memory.h
index 739e488ca3..2c9ec1b7c3 100644
--- a/include/hip/hcc_detail/hip_memory.h
+++ b/include/hip/hcc_detail/hip_memory.h
@@ -27,6 +27,8 @@ THE SOFTWARE.
 // HIP heap is implemented as a global array with fixed size. Users may define
 // __HIP_SIZE_OF_PAGE and __HIP_NUM_PAGES to have a larger heap.
 
+#if __HCC__ || __HIP__
+
 // Size of page in bytes.
 #ifndef __HIP_SIZE_OF_PAGE
 #define __HIP_SIZE_OF_PAGE 64
@@ -39,11 +41,8 @@ THE SOFTWARE.
 
 #define __HIP_SIZE_OF_HEAP (__HIP_NUM_PAGES * __HIP_SIZE_OF_PAGE)
 
-#if __HCC__ || __HIP__
-
-__attribute__((weak)) __device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP];
-__attribute__((weak)) __device__
-    uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
+extern __device__ char __hip_device_heap[];
+extern __device__ uint32_t __hip_device_page_flag[];
 
 extern "C" inline __device__ void* __hip_malloc(size_t size) {
     char* heap = (char*)__hip_device_heap;
diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index 93ac527826..4ea5b24f43 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -28,6 +28,9 @@ THE SOFTWARE.
 #include "hip_hcc_internal.h"
 #include "trace_helper.h"
 
+__device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP];
+__device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
+
 // Internal HIP APIS:
 namespace hip_internal {