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: 9ba3948388]
2017-06-12 13:31:09 -04:00
parent 4b1593bc58
commit 5fb0f84b9a
13 changed files with 273 additions and 34 deletions
@@ -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);
@@ -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<Device*>& devices() { return *devices_; }
+
+  // P2P devices that are accessible from the current device
+  std::vector<cl_device_id> p2pDevices_;
+
 protected:
  //! Enable the specified extension
  char* getExtensionString();
@@ -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);
@@ -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);
@@ -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;
@@ -48,6 +48,8 @@ amd::Device::Compiler* NullDevice::compilerHandle_;
 bool roc::Device::isHsaInitialized_ = false;
 hsa_agent_t roc::Device::cpu_agent_ = {0};
 std::vector<hsa_agent_t> roc::Device::gpu_agents_;
+amd::Monitor* roc::Device::p2p_stage_ops_ = nullptr;
+std::vector<Memory*> 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<Device*>(device1)->p2pAgents()) {
+      // Find cl_device_id associated with the current agent
+      for (auto device2: Device::devices()) {
+        if (agent.handle == static_cast<Device*>(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;
 }

@@ -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<hsa_agent_t>& 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<Memory*>& 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<amd::Memory*>* mapCache_;  //!< Map cache info structure

+  static amd::Monitor* p2p_stage_ops_;   //!< Lock to serialise cache for the P2P resources
+  static std::vector<Memory*> p2p_stages_; //!< Staging resources
+
  bool populateOCLDeviceConstants();
  static bool isHsaInitialized_;
  static hsa_agent_t cpu_agent_;
  static std::vector<hsa_agent_t> gpu_agents_;
+  std::vector<hsa_agent_t> p2p_agents_;  //!< List of P2P agents available for this device
  MesaInterop mesa_;
  hsa_agent_t _bkendDevice;
  hsa_profile_t agent_profile_;
@@ -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;
  }
@@ -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)
@@ -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<roc::Memory*>(
+    cmd.source().getDeviceMemory(*cmd.source().getContext().devices()[0]));
+  Memory* dstDevMem = static_cast<roc::Memory*>(
+    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
@@ -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) {}

@@ -558,4 +558,30 @@ bool TransferBufferFileCommand::validateMemory() {
  return true;
 }

+bool CopyMemoryP2PCommand::validateMemory() {
+  if (queue()->device().info().type_ & CL_DEVICE_TYPE_GPU) {
+    const std::vector<Device*>& 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<Device*>& 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
@@ -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();
+};
+
 /*! @}
 *  @}
 */