diff --git a/projects/clr/hipamd/src/hip_hmm.cpp b/projects/clr/hipamd/src/hip_hmm.cpp index 8a3cf7da16..652e4158c8 100644 --- a/projects/clr/hipamd/src/hip_hmm.cpp +++ b/projects/clr/hipamd/src/hip_hmm.cpp @@ -226,6 +226,9 @@ hipError_t ihipMallocManaged(void** ptr, size_t size, unsigned int align) { } size_t offset = 0; //this is ignored amd::Memory* memObj = getMemoryObject(*ptr, offset); + if (memObj == nullptr) { + return hipErrorMemoryAllocation; + } //saves the current device id so that it can be accessed later memObj->getUserData().deviceId = hip::getCurrentDevice()->deviceId(); diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 4be5e43e1c..75da3f0beb 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -43,13 +43,19 @@ amd::Memory* getMemoryObject(const void* ptr, size_t& offset, size_t size) { } else { //SVM ptr or device ptr mapped from host - const void *devPtr = reinterpret_cast - (memObj->getDeviceMemory(*memObj->getContext().devices()[0])->virtualAddress()); - if (devPtr != nullptr) { - offset = reinterpret_cast(ptr) - reinterpret_cast(devPtr); + device::Memory* devMem = + (memObj->getDeviceMemory(*memObj->getContext().devices()[0])); + if (devMem == nullptr) { + return nullptr; } else { - ShouldNotReachHere(); + const void* devPtr = reinterpret_cast((devMem)->virtualAddress()); + if (devPtr != nullptr) { + offset = reinterpret_cast(ptr) - reinterpret_cast(devPtr); + } + else { + return nullptr; + } } } } else {