Added hipHostAlloc with hipHostAllocMapped flag

[ROCm/clr commit: 411154f93f]

projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h

@@ -56,6 +56,11 @@ extern "C" {
 #define hipEventInterprocess  0x4  ///< Event can support IPC.  @warning - not supported in HIP.
 
 
+#define hipHostAllocDefault 0x0
+#define hipHostAllocPortable 0x1
+#define hipHostAllocMapped 0x2
+#define hipHostAllocWriteCombined 0x4
+
 /**
  * @warning On AMD devices and recent Nvidia devices, these hints and controls are ignored.
  */
@@ -662,6 +667,25 @@ hipError_t hipMalloc(void** ptr, size_t size) ;
  */
 hipError_t hipMallocHost(void** ptr, size_t size) ;
 
+/**
+ *  @brief Allocate device accessible page locked host memory 
+ *
+ *  @param[out]  ptr Pointer to the allocated host pinned memory
+ *  @param[in] size Requested memory size
+ *  @param[in] flags Type of host memory allocation
+ *  @return Error code
+ */
+hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags) ;
+
+/**
+ *  @brief Get Device pointer from Host Pointer allocated through hipHostAlloc
+ *
+ *  @param[out]  dstPtr Device Pointer mapped to passed host pointer
+ *  @param[in] hstPtr Host Pointer allocated through hipHostAlloc
+ *  @param[in] size Requested memory size
+ *  @return Error code
+ */
+hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, size_t size) ;
 
 /**
  *  @brief Free memory allocated by the hcc hip memory allocation API.

projects/clr/hipamd/src/hip_hcc.cpp

@@ -2000,6 +2000,73 @@ hipError_t hipMallocHost(void** ptr, size_t sizeBytes)
     return ihipLogStatus(hip_status);
 }
 
+hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags){
+	std::call_once(hip_initialized, ihipInit);
+
+	hipError_t hip_status = hipSuccess;
+
+	auto device = ihipGetTlsDefaultDevice();
+
+	if(device){
+		if(flags | hipHostAllocDefault){
+		const unsigned am_flags = amHostPinned;
+
+		*ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags);
+		if(sizeBytes && (*ptr == NULL)){
+			hip_status = hipErrorMemoryAllocation;
+		}else{
+#if USE_AM_TRACKER
+			hc::am_memtracker_update(*ptr, device->_device_index, 0);
+#endif
+		}
+		tprintf(TRACE_MEM, " %s: pinned ptr=%p\n", __func__, *ptr);
+		}
+		if(flags | hipHostAllocMapped && device->_props.canMapHostMemory == 1){
+		const unsigned am_flags = amHostPinned;
+
+		*ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags);
+		if(sizeBytes && (*ptr == NULL)){
+			hip_status = hipErrorMemoryAllocation;
+		}else{
+#if USE_AM_TRACKER
+			hc::am_memtracker_update(*ptr, device->_device_index, 0);
+			void *srcPtr;
+			hsa_status_t hsa_status = hsa_amd_memory_lock((*ptr), sizeBytes, &device->_hsa_agent, 1, &srcPtr);
+			assert(hsa_status == HSA_STATUS_SUCCESS);
+			hc::am_memtracker_add(srcPtr, sizeBytes, device->_acc, false);
+#endif
+		}
+		tprintf(TRACE_MEM, " %s: pinned ptr=%p\n", __func__, *ptr);
+		}
+	}
+	return ihipLogStatus(hip_status);
+}
+
+hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, size_t size){
+	std::call_once(hip_initialized, ihipInit);
+
+	hipError_t hip_status = hipSuccess;
+
+	if(hstPtr == NULL){
+		hip_status = hipErrorInvalidValue;
+	}else{
+
+#if USE_AM_TRACKER
+	hc::accelerator acc;
+	hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0);
+	am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hstPtr);
+	if(status == AM_SUCCESS){
+		*devPtr = amPointerInfo._devicePointer;
+		if(devPtr == NULL){
+			hip_status = hipErrorMemoryAllocation;
+	}
+	}
+#endif
+	tprintf(TRACE_MEM, " %s: pinned ptr=%p\n", __func__, *devPtr);
+	}
+	return ihipLogStatus(hip_status);
+}
+
 //---
 hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind)
 {