P4 to Git Change 1775995 by gandryey@gera-w8 on 2019/04/29 13:46:53

SWDEV-79445 - OCL generic changes and code clean-up
	- Enable P2P extension for PAL path, currently it's staging copy only
	- Fix P2P staging copy

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#55 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/device.cpp#243 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/device.hpp#336 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.cpp#130 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palsettings.cpp#77 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.cpp#131 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.hpp#59 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.cpp#124 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.hpp#36 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#74 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/command.cpp#92 edit

hipamd/api/hip/hip_memory.cpp

@@ -110,12 +110,6 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
         ((dstMemory->getMemFlags() & hostMem) == 0)) {
       amd::Device* queueDevice = &queue.device();
       if (queueDevice != srcMemory->getContext().devices()[0]) {
-        void* staging = nullptr;
-        ihipMalloc(&staging, sizeBytes, CL_MEM_SVM_FINE_GRAIN_BUFFER);
-        ihipMemcpy(staging, src, sizeBytes, hipMemcpyDeviceToHost, *hip::getNullStream(srcMemory->getContext()));
-        ihipMemcpy(dst, staging, sizeBytes, hipMemcpyHostToDevice, queue, isAsync);
-        hipFree(staging);
-#if 0
         amd::Coord3D srcOffset(sOffset, 0, 0);
         amd::Coord3D dstOffset(dOffset, 0, 0);
         amd::Coord3D copySize(sizeBytes, 1, 1);
@@ -126,16 +120,9 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
           command->awaitCompletion();
         }
         command->release();
-#endif
         return hipSuccess;
       }
       if (queueDevice != dstMemory->getContext().devices()[0]) {
-        void* staging = nullptr;
-        ihipMalloc(&staging, sizeBytes, CL_MEM_SVM_FINE_GRAIN_BUFFER);
-        ihipMemcpy(staging, src, sizeBytes, hipMemcpyDeviceToHost, queue);
-        ihipMemcpy(dst, staging, sizeBytes, hipMemcpyHostToDevice, *hip::getNullStream(dstMemory->getContext()), isAsync);
-        hipFree(staging);
-#if 0
         amd::Coord3D srcOffset(sOffset, 0, 0);
         amd::Coord3D dstOffset(dOffset, 0, 0);
         amd::Coord3D copySize(sizeBytes, 1, 1);
@@ -146,7 +133,6 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
           command->awaitCompletion();
         }
         command->release();
-#endif
         return hipSuccess;
       }
     }