From c5afd5d41230366b7eae35d3ce515baf1b793416 Mon Sep 17 00:00:00 2001 From: German Andryeyev Date: Thu, 11 Jun 2020 14:26:15 -0400 Subject: [PATCH] Initial HMM support - Expose ROCclr interfaces for HIP usage - ROCr interfaces aren't available in staging, thus control the build with AMD_HMM_SUPPORT define Change-Id: Iadc2bcc230e78d3b0dc22b235189c8cc80843446 --- rocclr/device/device.hpp | 47 +++++++- rocclr/device/rocm/CMakeLists.txt | 5 + rocclr/device/rocm/rocdevice.cpp | 187 +++++++++++++++++++++++++++++ rocclr/device/rocm/rocdevice.hpp | 17 ++- rocclr/device/rocm/rocmemory.cpp | 53 ++++++-- rocclr/device/rocm/rocsettings.cpp | 3 + rocclr/device/rocm/rocsettings.hpp | 9 ++ rocclr/device/rocm/rocvirtual.cpp | 26 ++++ rocclr/device/rocm/rocvirtual.hpp | 1 + rocclr/platform/command.cpp | 10 +- rocclr/platform/command.hpp | 24 +++- rocclr/utils/flags.hpp | 2 + 12 files changed, 368 insertions(+), 16 deletions(-) diff --git a/rocclr/device/device.hpp b/rocclr/device/device.hpp index aa3b695404..ef5fc98433 100755 --- a/rocclr/device/device.hpp +++ b/rocclr/device/device.hpp @@ -76,6 +76,7 @@ class SvmCopyMemoryCommand; class SvmFillMemoryCommand; class SvmMapMemoryCommand; class SvmUnmapMemoryCommand; +class SvmPrefetchAsyncCommand; class TransferBufferFileCommand; class HwDebugManager; class Device; @@ -86,6 +87,30 @@ namespace option { class Options; } // namespace option +//! @note: the defines match hip values +enum MemoryAdvice : uint32_t { + SetReadMostly = 1, ///< Data will mostly be read and only occassionally be written to + UnsetReadMostly = 2, ///< Undo the effect of hipMemAdviseSetReadMostly + SetPreferredLocation = 3, ///< Set the preferred location for the data as the specified device + UnsetPreferredLocation = 4, ///< Clear the preferred location for the data + SetAccessedBy = 5, ///< Data will be accessed by the specified device, + ///< so prevent page faults as much as possible + UnsetAccessedBy = 6 ///< Let the Unified Memory subsystem decide on + ///< the page faulting policy for the specified device +}; + +enum MemRangeAttribute : uint32_t { + ReadMostly = 1, ///< Whether the range will mostly be read and only + ///< occassionally be written to + PreferredLocation = 2, ///< The preferred location of the range + AccessedBy = 3, ///< Memory range has hipMemAdviseSetAccessedBy + ///< set for specified device + LastPrefetchLocation = 4, ///< The last location to which the range was prefetched +}; + +constexpr int CpuDeviceId = static_cast(-1); +constexpr int InvalidDeviceId = static_cast(-2); + } // namespace amd enum OclExtensions { @@ -1111,7 +1136,9 @@ class VirtualDevice : public amd::HeapObject { virtual void submitTransferBufferFromFile(amd::TransferBufferFileCommand& cmd) { ShouldNotReachHere(); } - + virtual void submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) { + ShouldNotReachHere(); + } //! Get the blit manager object device::BlitManager& blitMgr() const { return *blitMgr_; } @@ -1341,6 +1368,24 @@ class Device : public RuntimeObject { */ virtual void svmFree(void* ptr) const = 0; + /** + * @return True if the device successfully applied the SVM attributes in HMM for device memory + */ + virtual bool SetSvmAttributes(const void* dev_ptr, size_t count, + amd::MemoryAdvice advice, bool first_alloc = false) const { + ShouldNotCallThis(); + return false; + } + + /** + * @return True if the device successfully retrieved the SVM attributes from HMM for device memory + */ + virtual bool GetSvmAttributes(void** data, size_t* data_sizes, int* attributes, + size_t num_attributes, const void* dev_ptr, size_t count) const { + ShouldNotCallThis(); + return false; + } + //! Validate kernel virtual bool validateKernel(const amd::Kernel& kernel, const device::VirtualDevice* vdev, diff --git a/rocclr/device/rocm/CMakeLists.txt b/rocclr/device/rocm/CMakeLists.txt index 455fd7cc6c..3c3301b0f8 100644 --- a/rocclr/device/rocm/CMakeLists.txt +++ b/rocclr/device/rocm/CMakeLists.txt @@ -36,6 +36,11 @@ target_include_directories(oclrocm ${ROCM_OCL_INCLUDES} ${ROCR_INCLUDES}) +option(BUILD_HMM "Build HMM support" OFF) +if (BUILD_HMM) + target_compile_definitions(oclrocm + PRIVATE AMD_HMM_SUPPORT) +endif() if(USE_COMGR_LIBRARY) target_compile_definitions(oclrocm diff --git a/rocclr/device/rocm/rocdevice.cpp b/rocclr/device/rocm/rocdevice.cpp index 1ee1504983..acdb82b475 100644 --- a/rocclr/device/rocm/rocdevice.cpp +++ b/rocclr/device/rocm/rocdevice.cpp @@ -178,6 +178,7 @@ Device::Device(hsa_agent_t bkendDevice) system_coarse_segment_.handle = 0; gpuvm_segment_.handle = 0; gpu_fine_grained_segment_.handle = 0; + prefetch_signal_.handle = 0; } void Device::setupCpuAgent() { @@ -786,6 +787,11 @@ bool Device::create(bool sramEccEnabled) { } } + // Create signal for HMM prefetch operation on device + if (HSA_STATUS_SUCCESS != hsa_signal_create(InitSignalValue, 0, nullptr, &prefetch_signal_)) { + return false; + } + return true; } @@ -1904,6 +1910,187 @@ void* Device::svmAlloc(amd::Context& context, size_t size, size_t alignment, cl_ return svmPtr; } +// ================================================================================================ +bool Device::SetSvmAttributes(const void* dev_ptr, size_t count, + amd::MemoryAdvice advice, bool first_alloc) const { + if ((settings().hmmFlags_ & Settings::Hmm::EnableSvmTracking) && !first_alloc) { + amd::Memory* svmMem = svmMem = amd::MemObjMap::FindMemObj(dev_ptr); + if (nullptr == svmMem) { + LogPrintfError("SetSvmAttributes received unknown memory for update: %p!", dev_ptr); + return false; + } + } +#if AMD_HMM_SUPPORT + std::vector attr; + if (first_alloc) { + attr.push_back({HSA_AMD_SVM_ATTRIB_GLOBAL_FLAG, HSA_AMD_SVM_GLOBAL_FLAG_COARSE_GRAINED}); + } + + switch (advice) { + case amd::MemoryAdvice::SetReadMostly: + attr.push_back({HSA_AMD_SVM_ATTRIB_READ_ONLY, true}); + break; + case amd::MemoryAdvice::UnsetReadMostly: + attr.push_back({HSA_AMD_SVM_ATTRIB_READ_ONLY, false}); + break; + case amd::MemoryAdvice::SetPreferredLocation: + attr.push_back({HSA_AMD_SVM_ATTRIB_PREFERRED_LOCATION, getBackendDevice().handle}); + break; + case amd::MemoryAdvice::UnsetPreferredLocation: + // Note: The current behavior doesn't match hip spec precisely + attr.push_back({HSA_AMD_SVM_ATTRIB_PREFERRED_LOCATION, getCpuAgent().handle}); + break; + case amd::MemoryAdvice::SetAccessedBy: + attr.push_back({HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE, getBackendDevice().handle}); + break; + case amd::MemoryAdvice::UnsetAccessedBy: + // @note: The current behavior doesn't match hip spec precisely + attr.push_back({HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE, getCpuAgent().handle}); + break; + default: + return false; + break; + } + + hsa_status_t status = hsa_amd_svm_attributes_set(const_cast(dev_ptr), count, + attr.data(), attr.size()); + if (status != HSA_STATUS_SUCCESS) { + LogError("hsa_amd_svm_attributes_set() failed"); + return false; + } +#endif // AMD_HMM_SUPPORT + return true; +} + +// ================================================================================================ +bool Device::GetSvmAttributes(void** data, size_t* data_sizes, int* attributes, + size_t num_attributes, const void* dev_ptr, size_t count) const { + if (settings().hmmFlags_ & Settings::Hmm::EnableSvmTracking) { + amd::Memory* svmMem = svmMem = amd::MemObjMap::FindMemObj(dev_ptr); + if (nullptr == svmMem) { + LogPrintfError("GetSvmAttributes received unknown memory %p for state!", dev_ptr); + return false; + } + } +#if AMD_HMM_SUPPORT + std::vector attr; + + for (int i = 0; i < num_attributes; ++i) { + switch (attributes[i]) { + case amd::MemRangeAttribute::ReadMostly: + attr.push_back({HSA_AMD_SVM_ATTRIB_READ_ONLY, 0}); + break; + case amd::MemRangeAttribute::PreferredLocation: + attr.push_back({HSA_AMD_SVM_ATTRIB_PREFERRED_LOCATION, 0}); + break; + case amd::MemRangeAttribute::AccessedBy: + attr.push_back({HSA_AMD_SVM_ATTRIB_ACCESS_QUERY, 0}); + break; + case amd::MemRangeAttribute::LastPrefetchLocation: + attr.push_back({HSA_AMD_SVM_ATTRIB_PREFETCH_LOCATION, 0}); + break; + default: + return false; + break; + } + } + + hsa_status_t status = hsa_amd_svm_attributes_get(const_cast(dev_ptr), count, + attr.data(), attr.size()); + if (status != HSA_STATUS_SUCCESS) { + LogError("hsa_amd_svm_attributes_get() failed"); + return false; + } + + uint32_t idx = 0; + for (auto& it : attr) { + switch (it.attribute) { + case HSA_AMD_SVM_ATTRIB_READ_ONLY: + if (data_sizes[idx] != sizeof(uint32_t)) { + return false; + } + // Cast ROCr value into the hip format + *reinterpret_cast(data[idx]) = static_cast(it.value); + break; + // The logic should be identical for the both queries + case HSA_AMD_SVM_ATTRIB_PREFERRED_LOCATION: + case HSA_AMD_SVM_ATTRIB_PREFETCH_LOCATION: + if (data_sizes[idx] != sizeof(uint32_t)) { + return false; + } + *reinterpret_cast(data[idx]) = static_cast(amd::InvalidDeviceId); + // Find device agent returned by ROCr + for (auto& device : devices()) { + if (static_cast(device)->getBackendDevice().handle == it.value) { + *reinterpret_cast(data[idx]) = static_cast(device->index()); + } + } + // Find CPU agent returned by ROCr + for (auto& agent_info : getCpuAgents()) { + if (agent_info.agent.handle == it.value) { + *reinterpret_cast(data[idx]) = static_cast(amd::CpuDeviceId); + } + } + break; + case HSA_AMD_SVM_ATTRIB_ACCESS_QUERY: + // Make sure it's multiple of 4 + if (data_sizes[idx] % 4 != 0) { + return false; + } + // @note currently it's a nop + break; + default: + return false; + break; + } + // Find the next location in the query + ++idx; + } +#endif // AMD_HMM_SUPPORT + return true; +} + +// ================================================================================================ +bool Device::SvmAllocInit(void* memory, size_t size) const { + amd::MemoryAdvice advice = amd::MemoryAdvice::SetAccessedBy; + constexpr bool kFirstAlloc = true; + SetSvmAttributes(memory, size, advice, kFirstAlloc); + + if (settings().hmmFlags_ & Settings::Hmm::EnableSystemMemory) { + advice = amd::MemoryAdvice::UnsetPreferredLocation; + SetSvmAttributes(memory, size, advice); + } else { + advice = amd::MemoryAdvice::SetPreferredLocation; + SetSvmAttributes(memory, size, advice); + } + + if ((settings().hmmFlags_ & Settings::Hmm::EnableMallocPrefetch) == 0) { + return true; + } + +#if AMD_HMM_SUPPORT + // Initialize signal for the barrier + hsa_signal_store_relaxed(prefetch_signal_, InitSignalValue); + + // Initiate a prefetch command which should force memory update in HMM + hsa_status_t status = hsa_amd_svm_prefetch_async(memory, size, getBackendDevice(), + 0, nullptr, prefetch_signal_); + if (status != HSA_STATUS_SUCCESS) { + LogError("hsa_amd_svm_attributes_get() failed"); + return false; + } + + // Wait for the prefetch + if (hsa_signal_wait_acquire(prefetch_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), + HSA_WAIT_STATE_BLOCKED) != 0) { + LogError("Barrier packet submission failed"); + return false; + } +#endif // AMD_HMM_SUPPORT + return true; +} + +// ================================================================================================ void Device::svmFree(void* ptr) const { amd::Memory* svmMem = nullptr; svmMem = amd::MemObjMap::FindMemObj(ptr); diff --git a/rocclr/device/rocm/rocdevice.hpp b/rocclr/device/rocm/rocdevice.hpp index 54b49ef04f..039509f26d 100644 --- a/rocclr/device/rocm/rocdevice.hpp +++ b/rocclr/device/rocm/rocdevice.hpp @@ -79,7 +79,7 @@ class IProDevice; class Sampler : public device::Sampler { public: //! Constructor - Sampler(const Device& dev) : dev_(dev) {} + Sampler(const Device& dev) : dev_(dev) {} //! Default destructor for the device memory object virtual ~Sampler(); @@ -381,7 +381,13 @@ class Device : public NullDevice { virtual void svmFree(void* ptr) const; - virtual bool SetClockMode(const cl_set_device_clock_mode_input_amd setClockModeInput, cl_set_device_clock_mode_output_amd* pSetClockModeOutput); + virtual bool SetSvmAttributes(const void* dev_ptr, size_t count, + amd::MemoryAdvice advice, bool first_alloc = false) const; + virtual bool GetSvmAttributes(void** data, size_t* data_sizes, int* attributes, + size_t num_attributes, const void* dev_ptr, size_t count) const; + + virtual bool SetClockMode(const cl_set_device_clock_mode_input_amd setClockModeInput, + cl_set_device_clock_mode_output_amd* pSetClockModeOutput); //! Returns transfer engine object const device::BlitManager& xferMgr() const { return xferQueue()->blitMgr(); } @@ -464,7 +470,12 @@ class Device : public NullDevice { roc::Memory* getGpuMemory(amd::Memory* mem //!< Pointer to AMD memory object ) const; + //! Initialize memory in AMD HMM on the current device or keeps it in the host memory + bool SvmAllocInit(void* memory, size_t size) const; + private: + static const hsa_signal_value_t InitSignalValue = 1; + static hsa_ven_amd_loader_1_00_pfn_t amd_loader_ext_table; amd::Monitor* mapCacheOps_; //!< Lock to serialise cache for the map resources @@ -485,6 +496,8 @@ class Device : public NullDevice { hsa_amd_memory_pool_t system_coarse_segment_; hsa_amd_memory_pool_t gpuvm_segment_; hsa_amd_memory_pool_t gpu_fine_grained_segment_; + hsa_signal_t prefetch_signal_; //!< Prefetch signal, used to explicitly prefetch SVM on device + size_t gpuvm_segment_max_alloc_; size_t alloc_granularity_; static const bool offlineDevice_; diff --git a/rocclr/device/rocm/rocmemory.cpp b/rocclr/device/rocm/rocmemory.cpp index 6db9379004..a1d864f6ba 100755 --- a/rocclr/device/rocm/rocmemory.cpp +++ b/rocclr/device/rocm/rocmemory.cpp @@ -42,7 +42,7 @@ namespace roc { -/////////////////////////////////roc::Memory////////////////////////////// +// ======================================= roc::Memory ============================================ Memory::Memory(const roc::Device& dev, amd::Memory& owner) : device::Memory(owner), dev_(dev), @@ -620,8 +620,7 @@ void Memory::mgpuCacheWriteBack() { } } -/////////////////////////////////roc::Buffer////////////////////////////// - +// ==================================== roc::Buffer =============================================== Buffer::Buffer(const roc::Device& dev, amd::Memory& owner) : roc::Memory(dev, owner) {} Buffer::Buffer(const roc::Device& dev, size_t size) : roc::Memory(dev, size) {} @@ -634,6 +633,7 @@ Buffer::~Buffer() { } } +// ================================================================================================ void Buffer::destroy() { if (owner()->parent() != nullptr) { return; @@ -647,15 +647,24 @@ void Buffer::destroy() { cl_mem_flags memFlags = owner()->getMemFlags(); if (owner()->getSvmPtr() != nullptr) { - if (dev().forceFineGrain(owner()) || - dev().isFineGrainedSystem(true)) { + if (dev().forceFineGrain(owner()) || dev().isFineGrainedSystem(true)) { memFlags |= CL_MEM_SVM_FINE_GRAIN_BUFFER; } const bool isFineGrain = memFlags & CL_MEM_SVM_FINE_GRAIN_BUFFER; if (kind_ != MEMORY_KIND_PTRGIVEN) { if (isFineGrain) { - dev().hostFree(deviceMemory_, size()); + if (memFlags & CL_MEM_ALLOC_HOST_PTR) { + #if AMD_HMM_SUPPORT + // AMD HMM path. Destroy system memory + amd::Os::uncommitMemory(deviceMemory_, size()); + amd::Os::releaseMemory(deviceMemory_, size()); + #else + dev().hostFree(deviceMemory_, size());; + #endif // AMD_HMM_SUPPORT + } else { + dev().hostFree(deviceMemory_, size()); + } } else { dev().memFree(deviceMemory_, size()); } @@ -705,6 +714,7 @@ void Buffer::destroy() { } } +// ================================================================================================ bool Buffer::create() { if (owner() == nullptr) { deviceMemory_ = dev().hostAlloc(size(), 1, false); @@ -731,7 +741,19 @@ bool Buffer::create() { if (owner()->getSvmPtr() == reinterpret_cast(1)) { if (isFineGrain) { - if (memFlags & CL_MEM_SVM_ATOMICS) { + if (memFlags & CL_MEM_ALLOC_HOST_PTR) { +#if AMD_HMM_SUPPORT + // AMD HMM path. Just allocate system memory and KFD will manage it + deviceMemory_ = amd::Os::reserveMemory( + 0, size(), amd::Os::pageSize(), amd::Os::MEM_PROT_RW); + amd::Os::commitMemory(deviceMemory_, size(), amd::Os::MEM_PROT_RW); + // Currently HMM requires cirtain initial calls to mark sysmem allocation as + // GPU accessible or prefetch memory into GPU + dev().SvmAllocInit(deviceMemory_, size()); +#else + deviceMemory_ = dev().hostAlloc(size(), 1, false); +#endif // AMD_HMM_SUPPORT + } else if (memFlags & CL_MEM_SVM_ATOMICS) { deviceMemory_ = dev().hostAlloc(size(), 1, true); } else { @@ -745,6 +767,14 @@ bool Buffer::create() { } else { deviceMemory_ = owner()->getSvmPtr(); kind_ = MEMORY_KIND_PTRGIVEN; +#if AMD_HMM_SUPPORT + if (memFlags & CL_MEM_ALLOC_HOST_PTR) { + // Currently HMM requires cirtain initial calls to mark sysmem allocation as + // GPU accessible or prefetch memory into the current device + // @note: Skip any allocaiton here, since sysmem was allocated on another device. + dev().SvmAllocInit(deviceMemory_, size()); + } +#endif // AMD_HMM_SUPPORT } if (!isFineGrain && (owner()->parent() != nullptr) && @@ -870,9 +900,10 @@ bool Buffer::create() { if (owner()->getSvmPtr() != owner()->getHostMem()) { if (memFlags & (CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR)) { - hsa_amd_memory_pool_t pool = (memFlags & CL_MEM_SVM_ATOMICS)? dev().SystemSegment() : dev().SystemCoarseSegment(); - hsa_status_t status = hsa_amd_memory_lock_to_pool(owner()->getHostMem(), owner()->getSize(), nullptr, - 0, pool, 0, &deviceMemory_); + hsa_amd_memory_pool_t pool = (memFlags & CL_MEM_SVM_ATOMICS) ? + dev().SystemSegment() : dev().SystemCoarseSegment(); + hsa_status_t status = hsa_amd_memory_lock_to_pool(owner()->getHostMem(), + owner()->getSize(), nullptr, 0, pool, 0, &deviceMemory_); if (status != HSA_STATUS_SUCCESS) { DevLogPrintfError("Failed to lock memory to pool, failed with hsa_status: %d \n", status); deviceMemory_ = nullptr; @@ -887,7 +918,7 @@ bool Buffer::create() { return deviceMemory_ != nullptr; } -/////////////////////////////////roc::Image////////////////////////////// +// ======================================= roc::Image ============================================= typedef struct ChannelOrderMap { uint32_t cl_channel_order; hsa_ext_image_channel_order_t hsa_channel_order; diff --git a/rocclr/device/rocm/rocsettings.cpp b/rocclr/device/rocm/rocsettings.cpp index 8fa006d82b..2dec296390 100644 --- a/rocclr/device/rocm/rocsettings.cpp +++ b/rocclr/device/rocm/rocsettings.cpp @@ -86,6 +86,9 @@ Settings::Settings() { lcWavefrontSize64_ = true; imageBufferWar_ = false; + + hmmFlags_ = (!flagIsDefault(ROC_HMM_FLAGS)) ? ROC_HMM_FLAGS : + Hmm::EnableSystemMemory | Hmm::EnableMallocPrefetch; } bool Settings::create(bool fullProfile, int gfxipMajor, int gfxipMinor, bool coop_groups) { diff --git a/rocclr/device/rocm/rocsettings.hpp b/rocclr/device/rocm/rocsettings.hpp index 19fdccc777..d57a590f04 100644 --- a/rocclr/device/rocm/rocsettings.hpp +++ b/rocclr/device/rocm/rocsettings.hpp @@ -34,6 +34,13 @@ namespace roc { //! Device settings class Settings : public device::Settings { public: + enum Hmm : uint32_t { + EnableSystemMemory = 0x01, //!< Forces system memory preference by default + EnableMallocPrefetch = 0x02, //!< Skips default prefetch after allocation + EnableSvmTracking = 0x04, //!< Enables SW SVM tracking + EnableDebugSvm = 0x08 //!< Extra debug flag (reserved for runtime developers) + }; + union { struct { uint doublePrecision_ : 1; //!< Enables double precision support @@ -75,6 +82,8 @@ class Settings : public device::Settings { size_t sdmaCopyThreshold_; //!< Use SDMA to copy above this size + uint32_t hmmFlags_; //!< HMM functionality control flags + //! Default constructor Settings(); diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index 955f083805..8436d72dfc 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -1170,6 +1170,7 @@ void VirtualGPU::submitWriteMemory(amd::WriteMemoryCommand& cmd) { profilingEnd(cmd); } +// ================================================================================================ void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& cmd) { // Make sure VirtualGPU has an exclusive access to the resources amd::ScopedLock lock(execution()); @@ -1191,6 +1192,31 @@ void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& cmd) { profilingEnd(cmd); } +// ================================================================================================ +void VirtualGPU::submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) { +#if ROCR_HMM_SUPPORT + // Initialize signal for the barrier + hsa_signal_store_relaxed(barrier_signal_, InitSignalValue); + + // Find the requested agent for the transfer + hsa_agent_t agent = cmd.cpu_access() ? dev().cpu_agent() ? gpu_device(); + + // Initiate a prefetch command + hsa_amd_svm_prefetch_async(cmd.dev_prt(), cmd.count(), agent, 0, nullptr, barrier_signal_); + + // Wait for the prefetch + if (hsa_signal_wait_acquire(barrier_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), + HSA_WAIT_STATE_BLOCKED) != 0) { + LogError("Barrier packet submission failed"); + return false; + } + + // Add system scope, since the prefetch scope is unclear + addSystemScope(); +#endif // ROCR_HMM_SUPPORT +} + +// ================================================================================================ bool VirtualGPU::copyMemory(cl_command_type type, amd::Memory& srcMem, amd::Memory& dstMem, bool entire, const amd::Coord3D& srcOrigin, const amd::Coord3D& dstOrigin, const amd::Coord3D& size, diff --git a/rocclr/device/rocm/rocvirtual.hpp b/rocclr/device/rocm/rocvirtual.hpp index e3a9cfd879..6266553caf 100644 --- a/rocclr/device/rocm/rocvirtual.hpp +++ b/rocclr/device/rocm/rocvirtual.hpp @@ -208,6 +208,7 @@ class VirtualGPU : public device::VirtualDevice { void submitSvmFillMemory(amd::SvmFillMemoryCommand& cmd); void submitSvmMapMemory(amd::SvmMapMemoryCommand& cmd); void submitSvmUnmapMemory(amd::SvmUnmapMemoryCommand& cmd); + void submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd); // { roc OpenCL integration // Added these stub (no-ops) implementation of pure virtual methods, diff --git a/rocclr/platform/command.cpp b/rocclr/platform/command.cpp index 1e15aebead..6a361cbc8f 100644 --- a/rocclr/platform/command.cpp +++ b/rocclr/platform/command.cpp @@ -203,7 +203,6 @@ bool Event::awaitCompletion() { lock_.wait(); } } - ClPrint(LOG_DEBUG, LOG_WAIT, "event %p wait completed", this); } @@ -633,4 +632,13 @@ bool CopyMemoryP2PCommand::validateMemory() { return true; } +// ================================================================================================ +bool SvmPrefetchAsyncCommand::validateMemory() { + amd::Memory* svmMem = svmMem = amd::MemObjMap::FindMemObj(dev_ptr()); + if (nullptr == svmMem) { + LogPrintfError("SvmPrefetchAsync received unknown memory for prefetch: %p!", dev_ptr()); + return false; + } +} + } // namespace amd diff --git a/rocclr/platform/command.hpp b/rocclr/platform/command.hpp index 8fa381f1be..8342151b2f 100644 --- a/rocclr/platform/command.hpp +++ b/rocclr/platform/command.hpp @@ -1364,7 +1364,6 @@ class TransferBufferFileCommand : public OneMemoryArgCommand { * 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, @@ -1378,6 +1377,29 @@ class CopyMemoryP2PCommand : public CopyMemoryCommand { bool validateMemory(); }; +/*! \brief Prefetch command for SVM memory + * + * \details Prefetches SVM memory into the current device or CPU + */ +class SvmPrefetchAsyncCommand : public Command { + const void* dev_ptr_; //!< Device pointer to memory for prefetch + size_t count_; //!< the size for prefetch + bool cpu_access_; //!< Prefetch data into CPU location + + public: + SvmPrefetchAsyncCommand(HostQueue& queue, const EventWaitList& eventWaitList, + const void* dev_ptr, size_t count, bool cpu_access) + : Command(queue, 1, eventWaitList), dev_ptr_(dev_ptr), count_(count), cpu_access_(cpu_access) {} + + virtual void submit(device::VirtualDevice& device) { device.submitSvmPrefetchAsync(*this); } + + bool validateMemory(); + + const void* dev_ptr() const { return dev_ptr_; } + size_t count() const { return count_; } + size_t cpu_access() const { return cpu_access_; } +}; + /*! @} * @} */ diff --git a/rocclr/utils/flags.hpp b/rocclr/utils/flags.hpp index d630c413ee..da982e8892 100644 --- a/rocclr/utils/flags.hpp +++ b/rocclr/utils/flags.hpp @@ -56,6 +56,8 @@ debug(size_t, PARAMETERS_MIN_ALIGNMENT, 16, \ "Minimum alignment required for the abstract parameters stack") \ debug(size_t, MEMOBJ_BASE_ADDR_ALIGN, 4*Ki, \ "Alignment of the base address of any allocate memory object") \ +release(uint, ROC_HMM_FLAGS, 0, \ + "ROCm HMM configuration flags") \ release(cstring, GPU_DEVICE_ORDINAL, "", \ "Select the device ordinal (comma seperated list of available devices)") \ release(bool, REMOTE_ALLOC, false, \