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

@@ -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

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__
 

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) {

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 {

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"