From fa498d73d0f47a2aacac967b5233d400a6451217 Mon Sep 17 00:00:00 2001 From: German Andryeyev Date: Fri, 7 May 2021 14:22:19 -0400 Subject: [PATCH] SWDEV-240804 - HMM changes - Accept hipMemAttachHost flag - Don't allocate HMM memory if HMM is disabled in KFD Change-Id: I3d386eb0a109a19c16a87a728944a3a6669f643a [ROCm/hip commit: 31ae72a164992715eb216cf09d1a23ab07e9b812] --- projects/hip/rocclr/hip_hmm.cpp | 16 +++++++++++----- 1 file changed, 11 insertions(+), 5 deletions(-) diff --git a/projects/hip/rocclr/hip_hmm.cpp b/projects/hip/rocclr/hip_hmm.cpp index 6f5ac7225b..8829bb2753 100644 --- a/projects/hip/rocclr/hip_hmm.cpp +++ b/projects/hip/rocclr/hip_hmm.cpp @@ -59,7 +59,8 @@ static_assert(static_cast(hipMemRangeAttributeLastPrefetchLocation) == hipError_t hipMallocManaged(void** dev_ptr, size_t size, unsigned int flags) { HIP_INIT_API(hipMallocManaged, dev_ptr, size, flags); - if ((dev_ptr == nullptr) || (size == 0) || (flags != hipMemAttachGlobal)) { + if ((dev_ptr == nullptr) || (size == 0) || + ((flags != hipMemAttachGlobal) && (flags != hipMemAttachHost))) { HIP_RETURN(hipErrorInvalidValue); } @@ -187,19 +188,24 @@ hipError_t hipStreamAttachMemAsync(hipStream_t stream, hipDeviceptr_t* dev_ptr, // ================================================================================================ hipError_t ihipMallocManaged(void** ptr, size_t size, unsigned int align) { - if (size == 0) { + if (ptr == nullptr) { + return hipErrorInvalidValue; + } else if (size == 0) { *ptr = nullptr; return hipSuccess; - } else if (ptr == nullptr) { - return hipErrorInvalidValue; } assert((hip::host_device->asContext()!= nullptr) && "Current host context must be valid"); amd::Context& ctx = *hip::host_device->asContext(); const amd::Device& dev = *ctx.devices()[0]; + + if (!dev.info().hmmSupported_) { + return hipErrorInvalidValue; + } + // For now limit to the max allocation size on the device. - // The apps should be able to go over theCould you limit allocation in the future + // The apps should be able to go over the limit in the future if (dev.info().maxMemAllocSize_ < size) { return hipErrorMemoryAllocation; }