AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

P4 to Git Change 1766264 by cpaquot@cpaquot-ocl-lc-lnx on 2019/04/05 11:38:18

Browse Source 
SWDEV-145570 - [HIP] Use a context with all devices in system for host register
	hipHostRegister and hipMemcpy 0x10 and 0x20 fail in mGPU systems because
	we only register the memory on the current device. But in HIP, the registering
	needs to happen on all devices.

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_context.cpp#17 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#26 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#50 edit
This commit is contained in:
foreman
2019-04-05 11:58:25 -04:00
parent 22c104f84f
commit 8753616c97
3 changed files with 16 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files
api/hip/hip_memory.cpp
+7 -5
View File
@@ -62,11 +62,14 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
return hipErrorInvalidValue;
}
if (hip::getCurrentContext()->devices()[0]->info().maxMemAllocSize_ < sizeBytes) {
amd::Context* amdContext = ((flags & CL_MEM_SVM_FINE_GRAIN_BUFFER) != 0)?
hip::host_context : hip::getCurrentContext();
if (amdContext->devices()[0]->info().maxMemAllocSize_ < sizeBytes) {
return hipErrorMemoryAllocation;
}
*ptr = amd::SvmBuffer::malloc(*hip::getCurrentContext(), flags, sizeBytes, hip::getCurrentContext()->devices()[0]->info().memBaseAddrAlign_);
*ptr = amd::SvmBuffer::malloc(*amdContext, flags, sizeBytes, amdContext->devices()[0]->info().memBaseAddrAlign_);
if (*ptr == nullptr) {
return hipErrorOutOfMemory;
}
@@ -518,8 +521,7 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) {
hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) {
HIP_INIT_API(hostPtr, sizeBytes, flags);
if(hostPtr != nullptr) {
amd::Context *amdContext = hip::getCurrentContext();
amd::Memory* mem = new (*amdContext) amd::Buffer(*amdContext, CL_MEM_USE_HOST_PTR, sizeBytes);
amd::Memory* mem = new (*hip::host_context) amd::Buffer(*hip::host_context, CL_MEM_USE_HOST_PTR, sizeBytes);
if (!mem->create(hostPtr)) {
mem->release();
@@ -538,7 +540,7 @@ hipError_t hipHostUnregister(void* hostPtr) {
if (amd::SvmBuffer::malloced(hostPtr)) {
hip::syncStreams();
hip::getNullStream()->finish();
amd::SvmBuffer::free(*hip::getCurrentContext(), hostPtr);
amd::SvmBuffer::free(*hip::host_context, hostPtr);
HIP_RETURN(hipSuccess);
} else {
size_t offset = 0;