From 5fb0f84b9a86a4dd5060b059b0d5981e66b335ab Mon Sep 17 00:00:00 2001 From: foreman Date: Mon, 12 Jun 2017 13:31:09 -0400 Subject: [PATCH] P4 to Git Change 1421208 by gandryey@gera-w8 on 2017/06/12 13:15:22 SWDEV-124171 - adding support for p2p OCL in rocm stack - Add cl_amd_copy_buffer_p2p extension for P2P transfers. The extension adds a new API entry - clEnqueueCopyBufferP2PAMD() which allows to transfer CL buffers between different CL contexts on different GPUs. If P2P isn't possible, then double copy performed - Also the app can query the P2P support capabilities for the device. A list of P2P accessible devices can be returned for the current device http://ocltc.amd.com/reviews/r/12913/ Affected files ... ... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_context.cpp#54 edit ... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_device.cpp#62 edit ... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_p2p_amd.cpp#1 add ... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_p2p_amd.h#1 add ... //depot/stg/opencl/drivers/opencl/api/opencl/khronos/headers/opencl2.0/CL/cl_ext.h#29 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/cpu/cpuvirtual.hpp#14 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/device.hpp#287 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpuvirtual.hpp#141 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.hpp#26 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.cpp#19 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.cpp#54 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.hpp#22 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocmemory.cpp#24 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocsettings.cpp#19 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#39 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.hpp#12 edit ... //depot/stg/opencl/drivers/opencl/runtime/platform/command.cpp#79 edit ... //depot/stg/opencl/drivers/opencl/runtime/platform/command.hpp#84 edit [ROCm/clr commit: 9ba3948388551cf48c672a1876d417365428c78a] --- .../rocclr/runtime/device/cpu/cpuvirtual.hpp | 1 + projects/clr/rocclr/runtime/device/device.hpp | 17 +++- .../rocclr/runtime/device/gpu/gpuvirtual.hpp | 1 + .../rocclr/runtime/device/pal/palvirtual.hpp | 1 + .../rocclr/runtime/device/rocm/rocblit.cpp | 68 ++++++++------- .../rocclr/runtime/device/rocm/rocdevice.cpp | 67 +++++++++++++++ .../rocclr/runtime/device/rocm/rocdevice.hpp | 14 ++++ .../rocclr/runtime/device/rocm/rocmemory.cpp | 3 +- .../runtime/device/rocm/rocsettings.cpp | 2 +- .../rocclr/runtime/device/rocm/rocvirtual.cpp | 82 +++++++++++++++++++ .../rocclr/runtime/device/rocm/rocvirtual.hpp | 3 +- .../clr/rocclr/runtime/platform/command.cpp | 26 ++++++ .../clr/rocclr/runtime/platform/command.hpp | 22 +++++ 13 files changed, 273 insertions(+), 34 deletions(-) diff --git a/projects/clr/rocclr/runtime/device/cpu/cpuvirtual.hpp b/projects/clr/rocclr/runtime/device/cpu/cpuvirtual.hpp index 3bb2d3394f..2f30cee28e 100644 --- a/projects/clr/rocclr/runtime/device/cpu/cpuvirtual.hpp +++ b/projects/clr/rocclr/runtime/device/cpu/cpuvirtual.hpp @@ -35,6 +35,7 @@ class VirtualCPU : public device::VirtualDevice { virtual void submitReadMemory(amd::ReadMemoryCommand& command); virtual void submitWriteMemory(amd::WriteMemoryCommand& command); virtual void submitCopyMemory(amd::CopyMemoryCommand& command); + virtual void submitCopyMemoryP2P(amd::CopyMemoryP2PCommand& command) {} virtual void submitMapMemory(amd::MapMemoryCommand& command); virtual void submitUnmapMemory(amd::UnmapMemoryCommand& command); virtual void submitKernel(amd::NDRangeKernelCommand& command); diff --git a/projects/clr/rocclr/runtime/device/device.hpp b/projects/clr/rocclr/runtime/device/device.hpp index 9440a044da..0f53019d22 100644 --- a/projects/clr/rocclr/runtime/device/device.hpp +++ b/projects/clr/rocclr/runtime/device/device.hpp @@ -40,6 +40,7 @@ class ReadMemoryCommand; class WriteMemoryCommand; class FillMemoryCommand; class CopyMemoryCommand; +class CopyMemoryP2PCommand; class MapMemoryCommand; class UnmapMemoryCommand; class MigrateMemObjectsCommand; @@ -117,6 +118,7 @@ enum OclExtensions { ClKhrMipMapImageWrites, ClKhrIlProgram, ClAMDLiquidFlash, + ClAmdCopyBufferP2P, ClExtTotal }; @@ -158,6 +160,7 @@ static const char* OclExtensionsString[] = {"cl_khr_fp64 ", "cl_khr_mipmap_image_writes ", "", "cl_amd_liquid_flash ", + "cl_amd_copy_buffer_p2p ", NULL}; static constexpr int AmdVendor = 0x1002; @@ -770,12 +773,15 @@ class Memory : public amd::HeapObject { } } - //! Returns state of memory direct access flag + //! Returns the state of memory direct access flag bool isHostMemDirectAccess() const { return (flags_ & HostMemoryDirectAccess) ? true : false; } - //! Returns state of host memory registration flag + //! Returns the state of host memory registration flag bool isHostMemoryRegistered() const { return (flags_ & HostMemoryRegistered) ? true : false; } + //! Returns the state of CPU uncached access + bool isCpuUncached() const { return (flags_ & MemoryCpuUncached) ? true : false; } + protected: enum Flags { HostMemoryDirectAccess = 0x00000001, //!< GPU has direct access to the host memory @@ -783,6 +789,7 @@ class Memory : public amd::HeapObject { PinnedMemoryAlloced = 0x00000004, //!< An extra pinned resource was allocated SubMemoryObject = 0x00000008, //!< Memory is sub-memory HostMemoryRegistered = 0x00000010, //!< Host memory was registered + MemoryCpuUncached = 0x00000020 //!< Memory is uncached on CPU access(slow read) }; uint flags_; //!< Memory object flags @@ -1384,6 +1391,7 @@ class VirtualDevice : public amd::HeapObject { virtual void submitReadMemory(amd::ReadMemoryCommand& cmd) = 0; virtual void submitWriteMemory(amd::WriteMemoryCommand& cmd) = 0; virtual void submitCopyMemory(amd::CopyMemoryCommand& cmd) = 0; + virtual void submitCopyMemoryP2P(amd::CopyMemoryP2PCommand& cmd) = 0; virtual void submitMapMemory(amd::MapMemoryCommand& cmd) = 0; virtual void submitUnmapMemory(amd::UnmapMemoryCommand& cmd) = 0; virtual void submitKernel(amd::NDRangeKernelCommand& command) = 0; @@ -1692,6 +1700,11 @@ class Device : public RuntimeObject { //! Finds GPU memory from virtual address device::Memory* findMemoryFromVA(const void* ptr, size_t* offset) const; + static std::vector& devices() { return *devices_; } + + // P2P devices that are accessible from the current device + std::vector p2pDevices_; + protected: //! Enable the specified extension char* getExtensionString(); diff --git a/projects/clr/rocclr/runtime/device/gpu/gpuvirtual.hpp b/projects/clr/rocclr/runtime/device/gpu/gpuvirtual.hpp index b5d51dd6b0..b461491add 100644 --- a/projects/clr/rocclr/runtime/device/gpu/gpuvirtual.hpp +++ b/projects/clr/rocclr/runtime/device/gpu/gpuvirtual.hpp @@ -192,6 +192,7 @@ class VirtualGPU : public device::VirtualDevice, public CALGSLContext { void submitReadMemory(amd::ReadMemoryCommand& vcmd); void submitWriteMemory(amd::WriteMemoryCommand& vcmd); void submitCopyMemory(amd::CopyMemoryCommand& vcmd); + void submitCopyMemoryP2P(amd::CopyMemoryP2PCommand& vcmd) {} void submitMapMemory(amd::MapMemoryCommand& vcmd); void submitUnmapMemory(amd::UnmapMemoryCommand& vcmd); void submitKernel(amd::NDRangeKernelCommand& vcmd); diff --git a/projects/clr/rocclr/runtime/device/pal/palvirtual.hpp b/projects/clr/rocclr/runtime/device/pal/palvirtual.hpp index 77f18faf84..65ee2f5ad1 100644 --- a/projects/clr/rocclr/runtime/device/pal/palvirtual.hpp +++ b/projects/clr/rocclr/runtime/device/pal/palvirtual.hpp @@ -273,6 +273,7 @@ class VirtualGPU : public device::VirtualDevice { void submitReadMemory(amd::ReadMemoryCommand& vcmd); void submitWriteMemory(amd::WriteMemoryCommand& vcmd); void submitCopyMemory(amd::CopyMemoryCommand& vcmd); + void submitCopyMemoryP2P(amd::CopyMemoryP2PCommand& vcmd) {} void submitMapMemory(amd::MapMemoryCommand& vcmd); void submitUnmapMemory(amd::UnmapMemoryCommand& vcmd); void submitKernel(amd::NDRangeKernelCommand& vcmd); diff --git a/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp b/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp index f2fe4fef34..3065f987bc 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp @@ -49,7 +49,8 @@ bool DmaBlitManager::readBuffer(device::Memory& srcMemory, void* dstHost, const amd::Coord3D& origin, const amd::Coord3D& size, bool entire) const { // Use host copy if memory has direct access - if (setup_.disableReadBuffer_ || gpuMem(srcMemory).isHostMemDirectAccess()) { + if (setup_.disableReadBuffer_ || + (srcMemory.isHostMemDirectAccess() && !srcMemory.isCpuUncached())) { return HostBlitManager::readBuffer(srcMemory, dstHost, origin, size, entire); } else { size_t srcSize = size[0]; @@ -132,7 +133,8 @@ bool DmaBlitManager::readBufferRect(device::Memory& srcMemory, void* dstHost, const amd::BufferRect& bufRect, const amd::BufferRect& hostRect, const amd::Coord3D& size, bool entire) const { // Use host copy if memory has direct access - if (setup_.disableReadBufferRect_ || gpuMem(srcMemory).isHostMemDirectAccess()) { + if (setup_.disableReadBufferRect_ || + (srcMemory.isHostMemDirectAccess() && !srcMemory.isCpuUncached())) { return HostBlitManager::readBufferRect(srcMemory, dstHost, bufRect, hostRect, size, entire); } else { Memory& xferBuf = dev().xferRead().acquire(); @@ -195,7 +197,7 @@ bool DmaBlitManager::writeBuffer(const void* srcHost, device::Memory& dstMemory, const amd::Coord3D& origin, const amd::Coord3D& size, bool entire) const { // Use host copy if memory has direct access - if (setup_.disableWriteBuffer_ || gpuMem(dstMemory).isHostMemDirectAccess()) { + if (setup_.disableWriteBuffer_ || dstMemory.isHostMemDirectAccess()) { return HostBlitManager::writeBuffer(srcHost, dstMemory, origin, size, entire); } else { size_t dstSize = size[0]; @@ -330,8 +332,8 @@ bool DmaBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& dstMe const amd::Coord3D& srcOrigin, const amd::Coord3D& dstOrigin, const amd::Coord3D& size, bool entire) const { if (setup_.disableCopyBuffer_ || - (gpuMem(srcMemory).isHostMemDirectAccess() && (dev().agent_profile() != HSA_PROFILE_FULL) && - gpuMem(dstMemory).isHostMemDirectAccess())) { + (srcMemory.isHostMemDirectAccess() && !srcMemory.isCpuUncached() && + (dev().agent_profile() != HSA_PROFILE_FULL) && dstMemory.isHostMemDirectAccess())) { return HostBlitManager::copyBuffer(srcMemory, dstMemory, srcOrigin, dstOrigin, size); } else { return hsaCopy(gpuMem(srcMemory), gpuMem(dstMemory), srcOrigin, dstOrigin, size); @@ -344,7 +346,8 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d const amd::BufferRect& srcRect, const amd::BufferRect& dstRect, const amd::Coord3D& size, bool entire) const { if (setup_.disableCopyBufferRect_ || - (gpuMem(srcMemory).isHostMemDirectAccess() && gpuMem(dstMemory).isHostMemDirectAccess())) { + (srcMemory.isHostMemDirectAccess() && !srcMemory.isCpuUncached() && + dstMemory.isHostMemDirectAccess())) { return HostBlitManager::copyBufferRect(srcMemory, dstMemory, srcRect, dstRect, size, entire); } else { return false; @@ -502,11 +505,20 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory, return (status == HSA_STATUS_SUCCESS); } - // Detect the agents for memory allocations - const hsa_agent_t srcAgent = + hsa_agent_t srcAgent; + hsa_agent_t dstAgent; + + if (&srcMemory.dev() == &dstMemory.dev()) { + // Detect the agents for memory allocations + srcAgent = (srcMemory.isHostMemDirectAccess()) ? dev().getCpuAgent() : dev().getBackendDevice(); - const hsa_agent_t dstAgent = + dstAgent = (dstMemory.isHostMemDirectAccess()) ? dev().getCpuAgent() : dev().getBackendDevice(); + } + else { + srcAgent = srcMemory.dev().getBackendDevice(); + dstAgent = dstMemory.dev().getBackendDevice(); + } const hsa_signal_value_t kInitVal = 1; hsa_signal_store_relaxed(completion_signal_, kInitVal); @@ -797,7 +809,7 @@ bool KernelBlitManager::copyBufferToImage(device::Memory& srcMemory, device::Mem return result; } // Check if buffer is in system memory with direct access - else if (gpuMem(srcMemory).isHostMemDirectAccess() && + else if (srcMemory.isHostMemDirectAccess() && (((rowPitch == 0) && (slicePitch == 0)) || ((rowPitch == imgRowPitch) && ((slicePitch == 0) || (slicePitch == imgSlicePitch))))) { // First attempt to do this all with DMA, @@ -1000,7 +1012,7 @@ bool KernelBlitManager::copyImageToBuffer(device::Memory& srcMemory, device::Mem return result; } // Check if buffer is in system memory with direct access - else if (gpuMem(dstMemory).isHostMemDirectAccess() && + else if (dstMemory.isHostMemDirectAccess() && (((rowPitch == 0) && (slicePitch == 0)) || ((rowPitch == imgRowPitch) && ((slicePitch == 0) || (slicePitch == imgSlicePitch))))) { // First attempt to do this all with DMA, @@ -1323,9 +1335,8 @@ bool KernelBlitManager::readImage(device::Memory& srcMemory, void* dstHost, bool result = false; // Use host copy if memory has direct access - if (setup_.disableReadImage_ || (gpuMem(srcMemory).isHostMemDirectAccess())) { - result = - HostBlitManager::readImage(srcMemory, dstHost, origin, size, rowPitch, slicePitch, entire); + if (setup_.disableReadImage_ || (srcMemory.isHostMemDirectAccess() && !srcMemory.isCpuUncached())) { + result = HostBlitManager::readImage(srcMemory, dstHost, origin, size, rowPitch, slicePitch, entire); synchronize(); return result; } else { @@ -1369,9 +1380,8 @@ bool KernelBlitManager::writeImage(const void* srcHost, device::Memory& dstMemor bool result = false; // Use host copy if memory has direct access - if (setup_.disableWriteImage_ || gpuMem(dstMemory).isHostMemDirectAccess()) { - result = - HostBlitManager::writeImage(srcHost, dstMemory, origin, size, rowPitch, slicePitch, entire); + if (setup_.disableWriteImage_ || dstMemory.isHostMemDirectAccess()) { + result = HostBlitManager::writeImage(srcHost, dstMemory, origin, size, rowPitch, slicePitch, entire); synchronize(); return result; } else { @@ -1417,10 +1427,9 @@ bool KernelBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory bool rejected = false; // Fall into the ROC path for rejected transfers - if (setup_.disableCopyBufferRect_ || gpuMem(srcMemory).isHostMemDirectAccess() || - gpuMem(dstMemory).isHostMemDirectAccess()) { - result = - HostBlitManager::copyBufferRect(srcMemory, dstMemory, srcRectIn, dstRectIn, sizeIn, entire); + if (setup_.disableCopyBufferRect_ || + srcMemory.isHostMemDirectAccess() || dstMemory.isHostMemDirectAccess()) { + result = DmaBlitManager::copyBufferRect(srcMemory, dstMemory, srcRectIn, dstRectIn, sizeIn, entire); if (result) { synchronize(); @@ -1529,7 +1538,7 @@ bool KernelBlitManager::readBuffer(device::Memory& srcMemory, void* dstHost, amd::ScopedLock k(lockXferOps_); bool result = false; // Use host copy if memory has direct access - if (setup_.disableReadBuffer_ || (gpuMem(srcMemory).isHostMemDirectAccess())) { + if (setup_.disableReadBuffer_ || (srcMemory.isHostMemDirectAccess() && !srcMemory.isCpuUncached())) { result = HostBlitManager::readBuffer(srcMemory, dstHost, origin, size, entire); synchronize(); return result; @@ -1576,7 +1585,8 @@ bool KernelBlitManager::readBufferRect(device::Memory& srcMemory, void* dstHost, bool result = false; // Use host copy if memory has direct access - if (setup_.disableReadBufferRect_ || gpuMem(srcMemory).isHostMemDirectAccess()) { + if (setup_.disableReadBufferRect_ || + (srcMemory.isHostMemDirectAccess() && !srcMemory.isCpuUncached())) { result = HostBlitManager::readBufferRect(srcMemory, dstHost, bufRect, hostRect, size, entire); synchronize(); return result; @@ -1621,7 +1631,7 @@ bool KernelBlitManager::writeBuffer(const void* srcHost, device::Memory& dstMemo bool result = false; // Use host copy if memory has direct access - if (setup_.disableWriteBuffer_ || gpuMem(dstMemory).isHostMemDirectAccess()) { + if (setup_.disableWriteBuffer_ || dstMemory.isHostMemDirectAccess()) { result = HostBlitManager::writeBuffer(srcHost, dstMemory, origin, size, entire); synchronize(); return result; @@ -1669,7 +1679,7 @@ bool KernelBlitManager::writeBufferRect(const void* srcHost, device::Memory& dst bool result = false; // Use host copy if memory has direct access - if (setup_.disableWriteBufferRect_ || gpuMem(dstMemory).isHostMemDirectAccess()) { + if (setup_.disableWriteBufferRect_ || dstMemory.isHostMemDirectAccess()) { result = HostBlitManager::writeBufferRect(srcHost, dstMemory, hostRect, bufRect, size, entire); synchronize(); return result; @@ -1717,7 +1727,7 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern, bool result = false; // Use host fill if memory has direct access - if (setup_.disableFillBuffer_ || gpuMem(memory).isHostMemDirectAccess()) { + if (setup_.disableFillBuffer_ || memory.isHostMemDirectAccess()) { result = HostBlitManager::fillBuffer(memory, pattern, patternSize, origin, size, entire); synchronize(); return result; @@ -1775,9 +1785,9 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds const amd::Coord3D& sizeIn, bool entire) const { amd::ScopedLock k(lockXferOps_); bool result = false; - + bool p2p = (&gpuMem(srcMemory).dev() != &gpuMem(dstMemory).dev()); if (setup_.disableHwlCopyBuffer_ || - (!gpuMem(srcMemory).isHostMemDirectAccess() && !gpuMem(dstMemory).isHostMemDirectAccess())) { + (!srcMemory.isHostMemDirectAccess() && !dstMemory.isHostMemDirectAccess() && !p2p)) { uint blitType = BlitCopyBuffer; size_t dim = 1; size_t globalWorkOffset[3] = {0, 0, 0}; @@ -1867,7 +1877,7 @@ bool KernelBlitManager::fillImage(device::Memory& memory, const void* pattern, bool result = false; // Use host fill if memory has direct access - if (setup_.disableFillImage_ || gpuMem(memory).isHostMemDirectAccess()) { + if (setup_.disableFillImage_ || memory.isHostMemDirectAccess()) { result = HostBlitManager::fillImage(memory, pattern, origin, size, entire); synchronize(); return result; diff --git a/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp b/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp index 12fa741536..99303e165f 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp @@ -48,6 +48,8 @@ amd::Device::Compiler* NullDevice::compilerHandle_; bool roc::Device::isHsaInitialized_ = false; hsa_agent_t roc::Device::cpu_agent_ = {0}; std::vector roc::Device::gpu_agents_; +amd::Monitor* roc::Device::p2p_stage_ops_ = nullptr; +std::vector roc::Device::p2p_stages_; const bool roc::Device::offlineDevice_ = false; const bool roc::NullDevice::offlineDevice_ = true; @@ -146,6 +148,14 @@ Device::~Device() { delete mapCache_; delete mapCacheOps_; + delete p2p_stage_ops_; + p2p_stage_ops_ = nullptr; + + for (auto buf: p2p_stages_) { + delete buf; + } + p2p_stages_.clear(); + // Destroy temporary buffers for read/write delete xferRead_; delete xferWrite_; @@ -500,6 +510,20 @@ bool Device::init() { } } + // Loop through all available devices + for (auto device1: Device::devices()) { + // Find all agents that can have access to the current device + for (auto agent: static_cast(device1)->p2pAgents()) { + // Find cl_device_id associated with the current agent + for (auto device2: Device::devices()) { + if (agent.handle == static_cast(device2)->getBackendDevice().handle) { + // Device2 can have access to device1 + device2->p2pDevices_.push_back(as_cl(device1)); + } + } + } + } + return true; } @@ -605,6 +629,22 @@ bool Device::create() { // Use just 1 entry by default for the map cache mapCache_->push_back(nullptr); + if (p2p_stage_ops_ == nullptr) { + p2p_stage_ops_ = new amd::Monitor("P2P Staging Lock", true); + if (nullptr == p2p_stage_ops_) { + return false; + } + for (uint i = 0; i < 2; i++) { + Memory* buf = new Buffer(*this, kP2PStagingSize); + if ((buf != nullptr) && buf->create()) { + p2p_stages_.push_back(buf); + } else { + delete buf; + return false; + } + } + } + if (settings().stagedXferSize_ != 0) { // Initialize staged write buffers if (settings().stagedXferWrite_) { @@ -776,6 +816,24 @@ bool Device::populateOCLDeviceConstants() { assert(group_segment_.handle != 0); + for (auto agent: gpu_agents_) { + if (agent.handle != _bkendDevice.handle) { + hsa_status_t err; + // Can current GPU have access to another GPU memory pool + hsa_amd_memory_pool_access_t access; + err = hsa_amd_agent_memory_pool_get_info(agent, gpuvm_segment_, HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access); + if (err != HSA_STATUS_SUCCESS) { + continue; + } + + // Find accessible p2p agents - i.e != HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED + if (HSA_AMD_MEMORY_POOL_ACCESS_ALLOWED_BY_DEFAULT == access || + HSA_AMD_MEMORY_POOL_ACCESS_DISALLOWED_BY_DEFAULT == access) { + p2p_agents_.push_back(agent); + } + } + } + size_t group_segment_size = 0; if (HSA_STATUS_SUCCESS != hsa_amd_memory_pool_get_info(group_segment_, HSA_AMD_MEMORY_POOL_INFO_SIZE, @@ -1307,6 +1365,15 @@ void* Device::deviceLocalAlloc(size_t size) const { return nullptr; } + if (p2pAgents().size() > 0) { + stat = hsa_amd_agents_allow_access(p2pAgents().size(), p2pAgents().data(), nullptr, ptr); + if (stat != HSA_STATUS_SUCCESS) { + LogError("Allow p2p acces for memory allocation"); + memFree(ptr, size); + return nullptr; + } + } + return ptr; } diff --git a/projects/clr/rocclr/runtime/device/rocm/rocdevice.hpp b/projects/clr/rocclr/runtime/device/rocm/rocdevice.hpp index 7e7d6fd7f7..7e348c0b3a 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocdevice.hpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocdevice.hpp @@ -209,6 +209,7 @@ class NullDevice : public amd::Device { //! A HSA device ordinal (physical HSA device) class Device : public NullDevice { public: + static constexpr size_t kP2PStagingSize = 1* Mi; //! Transfer buffers class XferBuffers : public amd::HeapObject { public: @@ -391,16 +392,29 @@ class Device : public NullDevice { const IProDevice& iPro() const { return *pro_device_; } bool ProEna() const { return pro_ena_; } + // P2P agents avaialble for this device + const std::vector& p2pAgents() const { return p2p_agents_; } + + // Lock protect P2P staging operations + amd::Monitor* P2PStageOps() const { return p2p_stage_ops_; } + + // Lock protect P2P staging operations + const std::vector& P2PStages() const { return p2p_stages_; } + private: static hsa_ven_amd_loader_1_00_pfn_t amd_loader_ext_table; amd::Monitor* mapCacheOps_; //!< Lock to serialise cache for the map resources std::vector* mapCache_; //!< Map cache info structure + static amd::Monitor* p2p_stage_ops_; //!< Lock to serialise cache for the P2P resources + static std::vector p2p_stages_; //!< Staging resources + bool populateOCLDeviceConstants(); static bool isHsaInitialized_; static hsa_agent_t cpu_agent_; static std::vector gpu_agents_; + std::vector p2p_agents_; //!< List of P2P agents available for this device MesaInterop mesa_; hsa_agent_t _bkendDevice; hsa_profile_t agent_profile_; diff --git a/projects/clr/rocclr/runtime/device/rocm/rocmemory.cpp b/projects/clr/rocclr/runtime/device/rocm/rocmemory.cpp index 757cb7cec5..1bc5240156 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocmemory.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocmemory.cpp @@ -621,6 +621,7 @@ bool Buffer::create() { deviceMemory_ = parentBuffer->getDeviceMemory() + offset; flags_ |= parentBuffer->isHostMemDirectAccess() ? HostMemoryDirectAccess : 0; + flags_ |= parentBuffer->isCpuUncached() ? MemoryCpuUncached : 0; // Explicitly set the host memory location, // because the parent location could change after reallocation @@ -643,7 +644,7 @@ bool Buffer::create() { if (deviceMemory_ == nullptr) { return false; } - flags_ |= HostMemoryDirectAccess; + flags_ |= HostMemoryDirectAccess | MemoryCpuUncached; owner()->setHostMem(host_ptr); return true; } diff --git a/projects/clr/rocclr/runtime/device/rocm/rocsettings.cpp b/projects/clr/rocclr/runtime/device/rocm/rocsettings.cpp index 3d3192cbd2..c9cc8c27c6 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocsettings.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocsettings.cpp @@ -107,8 +107,8 @@ bool Settings::create(bool fullProfile, int gfxipVersion) { enableExtension(ClAmdFp64); #endif // !defined(WITH_LIGHTNING_COMPILER) enableExtension(ClKhrSubGroups); - enableExtension(ClKhrDepthImages); + enableExtension(ClAmdCopyBufferP2P); supportDepthsRGB_ = true; #if defined(WITH_LIGHTNING_COMPILER) diff --git a/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp index ec630a9568..181593d205 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp @@ -1024,6 +1024,88 @@ void VirtualGPU::submitCopyMemory(amd::CopyMemoryCommand& cmd) { profilingEnd(cmd); } +void VirtualGPU::submitCopyMemoryP2P(amd::CopyMemoryP2PCommand& cmd) { + // Wait on a kernel if one is outstanding + releaseGpuMemoryFence(); + + profilingBegin(cmd); + + Memory* srcDevMem = static_cast( + cmd.source().getDeviceMemory(*cmd.source().getContext().devices()[0])); + Memory* dstDevMem = static_cast( + cmd.destination().getDeviceMemory(*cmd.destination().getContext().devices()[0])); + + bool p2pAllowed = false; + // Loop through all available P2P devices for the destination buffer + for (auto agent: dstDevMem->dev().p2pAgents()) { + // Find the device, which is matching the current + if (agent.handle == dev().getBackendDevice().handle) { + p2pAllowed = true; + break; + } + } + // Synchronize source and destination memory + device::Memory::SyncFlags syncFlags; + syncFlags.skipEntire_ = cmd.isEntireMemory(); + amd::Coord3D size = cmd.size(); + + bool result = false; + switch (cmd.type()) { + case CL_COMMAND_COPY_BUFFER: { + amd::Coord3D srcOrigin(cmd.srcOrigin()[0]); + amd::Coord3D dstOrigin(cmd.dstOrigin()[0]); + + if (p2pAllowed) { + result = blitMgr().copyBuffer(*srcDevMem, *dstDevMem, srcOrigin, dstOrigin, + size, cmd.isEntireMemory()); + } + else { + size_t copy_size = Device::kP2PStagingSize; + size_t left_size = size[0]; + result = true; + do { + if (left_size <= copy_size) { + copy_size = left_size; + } + left_size -= copy_size; + amd::Coord3D stageOffset(0); + amd::Coord3D cpSize(copy_size); + + // Perform 2 step transfer with staging buffer + // todo: optimization can be done with double buffering if events tracking + // will be propagated outside of the device transfers object + result &= dev().xferMgr().copyBuffer(*srcDevMem, *(dev().P2PStages()[0]), srcOrigin, + stageOffset, cpSize, cmd.isEntireMemory()); + srcOrigin.c[0] += copy_size; + result &= dstDevMem->dev().xferMgr().copyBuffer(*dstDevMem->dev().P2PStages()[0], + *dstDevMem, stageOffset, dstOrigin, + copy_size, cmd.isEntireMemory()); + dstOrigin.c[0] += copy_size; + } while (left_size > 0); + } + break; + } + case CL_COMMAND_COPY_BUFFER_RECT: + case CL_COMMAND_COPY_IMAGE: + case CL_COMMAND_COPY_IMAGE_TO_BUFFER: + case CL_COMMAND_COPY_BUFFER_TO_IMAGE: + LogError("Unsupported P2P type!"); + break; + default: + ShouldNotReachHere(); + break; + } + + if (!result) { + LogError("submitCopyMemoryP2P failed!"); + cmd.setStatus(CL_OUT_OF_RESOURCES); + } + + cmd.destination().signalWrite(&dstDevMem->dev()); + + profilingEnd(cmd); +} + void VirtualGPU::submitSvmMapMemory(amd::SvmMapMemoryCommand& cmd) { // No fence is needed since this is a no-op: the // command will be completed only after all the diff --git a/projects/clr/rocclr/runtime/device/rocm/rocvirtual.hpp b/projects/clr/rocclr/runtime/device/rocm/rocvirtual.hpp index 74bb4f8cf7..85075552c3 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocvirtual.hpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocvirtual.hpp @@ -158,6 +158,7 @@ class VirtualGPU : public device::VirtualDevice { void submitReadMemory(amd::ReadMemoryCommand& cmd); void submitWriteMemory(amd::WriteMemoryCommand& cmd); void submitCopyMemory(amd::CopyMemoryCommand& cmd); + void submitCopyMemoryP2P(amd::CopyMemoryP2PCommand& cmd); void submitMapMemory(amd::MapMemoryCommand& cmd); void submitUnmapMemory(amd::UnmapMemoryCommand& cmd); void submitKernel(amd::NDRangeKernelCommand& cmd); @@ -181,7 +182,7 @@ class VirtualGPU : public device::VirtualDevice { // Added these stub (no-ops) implementation of pure virtual methods, // when integrating HSA and OpenCL branches. // TODO: After inegration, whoever is working on VirtualGPU should write - // actual implemention. + // actual implementation. virtual void submitSignal(amd::SignalCommand& cmd) {} virtual void submitMakeBuffersResident(amd::MakeBuffersResidentCommand& cmd) {} diff --git a/projects/clr/rocclr/runtime/platform/command.cpp b/projects/clr/rocclr/runtime/platform/command.cpp index efe02a6802..b7cdae1c60 100644 --- a/projects/clr/rocclr/runtime/platform/command.cpp +++ b/projects/clr/rocclr/runtime/platform/command.cpp @@ -558,4 +558,30 @@ bool TransferBufferFileCommand::validateMemory() { return true; } +bool CopyMemoryP2PCommand::validateMemory() { + if (queue()->device().info().type_ & CL_DEVICE_TYPE_GPU) { + const std::vector& devices = memory1_->getContext().devices(); + if (devices.size() != 1) { + LogError("Can't allocate memory object for P2P extension"); + return false; + } + device::Memory* mem = memory1_->getDeviceMemory(*devices[0]); + if (nullptr == mem) { + LogPrintfError("Can't allocate memory size - 0x%08X bytes!", memory1_->getSize()); + return false; + } + const std::vector& devices2 = memory2_->getContext().devices(); + if (devices2.size() != 1) { + LogError("Can't allocate memory object for P2P extension"); + return false; + } + mem = memory2_->getDeviceMemory(*devices2[0]); + if (nullptr == mem) { + LogPrintfError("Can't allocate memory size - 0x%08X bytes!", memory2_->getSize()); + return false; + } + } + return true; +} + } // namespace amd diff --git a/projects/clr/rocclr/runtime/platform/command.hpp b/projects/clr/rocclr/runtime/platform/command.hpp index 5cfa667cca..c4b7fd4a8e 100644 --- a/projects/clr/rocclr/runtime/platform/command.hpp +++ b/projects/clr/rocclr/runtime/platform/command.hpp @@ -1292,6 +1292,28 @@ class TransferBufferFileCommand : public OneMemoryArgCommand { bool validateMemory(); }; +/*! \brief A P2P copy memory command + * + * \details Used for buffers only. Backends are expected + * to handle any required translation. Buffers are treated + * as 1D structures so origin_[0] and size_[0] are + * equivalent to offset_ and count_ respectively. + */ + +class CopyMemoryP2PCommand : public CopyMemoryCommand { + public: + CopyMemoryP2PCommand(HostQueue& queue, cl_command_type cmdType, const EventWaitList& eventWaitList, + Memory& srcMemory, Memory& dstMemory, Coord3D srcOrigin, Coord3D dstOrigin, + Coord3D size) + : CopyMemoryCommand(queue, cmdType, eventWaitList, srcMemory, dstMemory, srcOrigin, dstOrigin, size) + { + } + + virtual void submit(device::VirtualDevice& device) { device.submitCopyMemoryP2P(*this); } + + bool validateMemory(); +}; + /*! @} * @} */