From 8d83e954572e756fbb9ed34467de31a207666407 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu@amd.com>
Date: Tue, 14 Apr 2020 06:37:14 -0400
Subject: [PATCH] Disable device side malloc (#2009)

* Disable device side malloc

Currently device side malloc is not working and takes excessive
device memory.

Disable it for now until a working malloc is implemented.

Change-Id: I1ad908c1c53a83752383b4be96688a848642c699
---
 include/hip/hcc_detail/hip_memory.h     |  2 +-
 include/hip/hcc_detail/hip_runtime.h    | 11 ++++++++++-
 src/hip_device.cpp                      |  5 +++--
 src/hip_memory.cpp                      |  3 +++
 tests/src/deviceLib/hipDeviceMalloc.cpp |  4 ++--
 5 files changed, 19 insertions(+), 6 deletions(-)

diff --git a/include/hip/hcc_detail/hip_memory.h b/include/hip/hcc_detail/hip_memory.h
index 866b9e879e..0c006143de 100644
--- a/include/hip/hcc_detail/hip_memory.h
+++ b/include/hip/hcc_detail/hip_memory.h
@@ -27,7 +27,7 @@ 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__
+#if (__HCC__ || __HIP__) && __HIP_ENABLE_DEVICE_MALLOC__
 
 // Size of page in bytes.
 #ifndef __HIP_SIZE_OF_PAGE
diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h
index 0707cc6899..582e0cdefa 100644
--- a/include/hip/hcc_detail/hip_runtime.h
+++ b/include/hip/hcc_detail/hip_runtime.h
@@ -44,6 +44,11 @@ THE SOFTWARE.
 #include <stddef.h>
 #endif  //__cplusplus
 
+// __hip_malloc is not working. Disable it by default.
+#ifndef __HIP_ENABLE_DEVICE_MALLOC__
+#define __HIP_ENABLE_DEVICE_MALLOC__ 0
+#endif
+
 #if __HCC_OR_HIP_CLANG__
 
 #if __HIP__
@@ -305,11 +310,15 @@ static constexpr Coordinates<hip_impl::WorkitemId> threadIdx{};
 
 #endif // defined __HCC__
 #if __HCC_OR_HIP_CLANG__
+#if __HIP_ENABLE_DEVICE_MALLOC__
 extern "C" __device__ void* __hip_malloc(size_t);
 extern "C" __device__ void* __hip_free(void* ptr);
-
 static inline __device__ void* malloc(size_t size) { return __hip_malloc(size); }
 static inline __device__ void* free(void* ptr) { return __hip_free(ptr); }
+#else
+static inline __device__ void* malloc(size_t size) { __builtin_trap(); return nullptr; }
+static inline __device__ void* free(void* ptr) { __builtin_trap(); return nullptr; }
+#endif
 
 #endif //__HCC_OR_HIP_CLANG__
 
diff --git a/src/hip_device.cpp b/src/hip_device.cpp
index e5797727ae..f7d6b3ac79 100644
--- a/src/hip_device.cpp
+++ b/src/hip_device.cpp
@@ -96,12 +96,13 @@ hipError_t hipDeviceGetLimit(size_t* pValue, hipLimit_t limit) {
     if (pValue == nullptr) {
         return ihipLogStatus(hipErrorInvalidValue);
     }
+#if __HIP_ENABLE_DEVICE_MALLOC__
     if (limit == hipLimitMallocHeapSize) {
         *pValue = (size_t)__HIP_SIZE_OF_HEAP;
         return ihipLogStatus(hipSuccess);
-    } else {
-        return ihipLogStatus(hipErrorUnsupportedLimit);
     }
+#endif
+    return ihipLogStatus(hipErrorUnsupportedLimit);
 }
 
 hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t cacheConfig) {
diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index de6bc63b20..82ecaea82a 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -19,6 +19,7 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
+
 #include <hc_am.hpp>
 #include "hsa/hsa.h"
 #include "hsa/hsa_ext_amd.h"
@@ -30,8 +31,10 @@ THE SOFTWARE.
 #include <functional>
 #include <fstream>
 
+#if __HIP_ENABLE_DEVICE_MALLOC__
 __device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP];
 __device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
+#endif
 
 // Internal HIP APIS:
 namespace hip_internal {
diff --git a/tests/src/deviceLib/hipDeviceMalloc.cpp b/tests/src/deviceLib/hipDeviceMalloc.cpp
index 4af7614a5c..aec891ed19 100644
--- a/tests/src/deviceLib/hipDeviceMalloc.cpp
+++ b/tests/src/deviceLib/hipDeviceMalloc.cpp
@@ -17,8 +17,8 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 /* HIT_START
- * BUILD: %t %s NVCC_OPTIONS -std=c++11
- * TEST: %t EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s NVCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM all
+ * TEST: %t EXCLUDE_HIP_PLATFORM all
  * HIT_END
  */
 #include "test_common.h"