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

[ROCm/hip commit: 8d83e95457]

projects/hip/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

projects/hip/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__