diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_blit_kernel.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_blit_kernel.h index f6f7b27e23..eabbdc768a 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_blit_kernel.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_blit_kernel.h @@ -51,7 +51,7 @@ namespace amd { class BlitKernel : public core::Blit { public: - explicit BlitKernel(); + explicit BlitKernel(core::Queue* queue); virtual ~BlitKernel() override; /// @brief Initialize a blit kernel object. @@ -175,12 +175,9 @@ class BlitKernel : public core::Blit { std::map kernels_; /// AQL queue for submitting the vector copy kernel. - hsa_queue_t* queue_; + core::Queue* queue_; uint32_t queue_bitmask_; - /// Index to track concurrent kernel launch. - volatile uint64_t cached_index_; - /// Pointer to the kernel argument buffer. KernelArgs* kernarg_async_; uint32_t kernarg_async_mask_; diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_agent.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_agent.h index 7b32464653..56ddeee704 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_agent.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_agent.h @@ -79,10 +79,11 @@ class GpuAgentInt : public core::Agent { // @retval HSA_STATUS_SUCCESS DMA queue initialization is successful. virtual void InitDma() = 0; - // @brief Initialize blit kernel object based on AQL queue. + // @brief Initialization hook invoked after tools library has loaded, + // to allow tools interception of interface functions. // - // @retval HSA_STATUS_SUCCESS blit kernel object initialization is successful. - virtual hsa_status_t InitBlitKernel() = 0; + // @retval HSA_STATUS_SUCCESS if initialization is successful. + virtual hsa_status_t PostToolsInit() = 0; // @brief Invoke the user provided callback for each region accessible by // this agent. @@ -184,7 +185,7 @@ class GpuAgent : public GpuAgentInt { void InitDma() override; // @brief Override from core::Agent. - hsa_status_t InitBlitKernel() override; + hsa_status_t PostToolsInit() override; uint16_t GetMicrocodeVersion() const; @@ -309,15 +310,18 @@ class GpuAgent : public GpuAgentInt { static const uint32_t minAqlSize_ = 0x1000; // 4KB min static const uint32_t maxAqlSize_ = 0x20000; // 8MB max + // @brief Create a queue through HSA API to allow tools to intercept. + core::Queue* CreateInterceptibleQueue(); + // @brief Create SDMA blit object. // // @retval NULL if SDMA blit creation and initialization failed. core::Blit* CreateBlitSdma(); - // @brief Create Kernel blit object. + // @brief Create Kernel blit object using provided compute queue. // // @retval NULL if Kernel blit creation and initialization failed. - core::Blit* CreateBlitKernel(); + core::Blit* CreateBlitKernel(core::Queue* queue); // @brief Invoke the user provided callback for every region in @p regions. // @@ -359,15 +363,19 @@ class GpuAgent : public GpuAgentInt { // @brief Default scratch size per work item. size_t scratch_per_thread_; - // @brief Blit object to handle memory copy from system to device memory. - core::Blit* blit_h2d_; + // @brief Blit interfaces for each data path. + enum BlitEnum { BlitHostToDev, BlitDevToHost, BlitDevToDev, BlitCount }; - // @brief Blit object to handle memory copy from device to system memory. - core::Blit* blit_d2h_; + core::Blit* blits_[BlitCount]; - // @brief Blit object to handle memory copy from device to device memory, and - // memory fill. - core::Blit* blit_d2d_; + // @brief AQL queues for cache management and blit compute usage. + enum QueueEnum { + QueueUtility, // Cache management and device to {host,device} blit compute + QueueBlitOnly, // Host to device blit + QueueCount + }; + + core::Queue* queues_[QueueCount]; // @brief Mutex to protect the update to coherency type. KernelMutex coherency_lock_; diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp index aff60f08ba..1db1f4455a 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp @@ -515,10 +515,9 @@ static int kCopyMisalignedUnroll = GetKernelSourceParam("kCopyMisalignedUnroll") static int kFillVecWidth = GetKernelSourceParam("kFillVecWidth"); static int kFillUnroll = GetKernelSourceParam("kFillUnroll"); -BlitKernel::BlitKernel() +BlitKernel::BlitKernel(core::Queue* queue) : core::Blit(), - queue_(NULL), - cached_index_(0), + queue_(queue), kernarg_async_(NULL), kernarg_async_mask_(0), kernarg_async_counter_(0), @@ -529,40 +528,18 @@ BlitKernel::BlitKernel() BlitKernel::~BlitKernel() {} hsa_status_t BlitKernel::Initialize(const core::Agent& agent) { - hsa_agent_t agent_handle = agent.public_handle(); + queue_bitmask_ = queue_->public_handle()->size - 1; - uint32_t features = 0; - hsa_status_t status = - HSA::hsa_agent_get_info(agent_handle, HSA_AGENT_INFO_FEATURE, &features); - if (status != HSA_STATUS_SUCCESS) { - return status; - } - - if ((features & HSA_AGENT_FEATURE_KERNEL_DISPATCH) == 0) { - return HSA_STATUS_ERROR; - } - - status = HSA::hsa_queue_create(agent_handle, 1024, HSA_QUEUE_TYPE_MULTI, NULL, - NULL, 0, 0, &queue_); - - if (HSA_STATUS_SUCCESS != status) { - return status; - } - - queue_bitmask_ = queue_->size - 1; - - cached_index_ = 0; - - status = HSA::hsa_signal_create(1, 0, NULL, &completion_signal_); + hsa_status_t status = HSA::hsa_signal_create(1, 0, NULL, &completion_signal_); if (HSA_STATUS_SUCCESS != status) { return status; } kernarg_async_ = reinterpret_cast( core::Runtime::runtime_singleton_->system_allocator()( - queue_->size * AlignUp(sizeof(KernelArgs), 16), 16)); + queue_->public_handle()->size * AlignUp(sizeof(KernelArgs), 16), 16)); - kernarg_async_mask_ = queue_->size - 1; + kernarg_async_mask_ = queue_->public_handle()->size - 1; // Obtain the number of compute units in the underlying agent. const GpuAgent& gpuAgent = static_cast(agent); @@ -598,10 +575,6 @@ hsa_status_t BlitKernel::Destroy(const core::Agent& agent) { kernel_pair.second.code_buf_size_); } - if (queue_ != NULL) { - HSA::hsa_queue_destroy(queue_); - } - if (kernarg_async_ != NULL) { core::Runtime::runtime_singleton_->system_deallocator()(kernarg_async_); } @@ -661,7 +634,8 @@ hsa_status_t BlitKernel::SubmitLinearCopyCommand( barrier_packet.header = HSA_PACKET_TYPE_INVALID; hsa_barrier_and_packet_t* queue_buffer = - reinterpret_cast(queue_->base_address); + reinterpret_cast( + queue_->public_handle()->base_address); const size_t dep_signal_count = dep_signals.size(); for (size_t i = 0; i < dep_signal_count; ++i) { @@ -803,26 +777,20 @@ hsa_status_t BlitKernel::SubmitLinearFillCommand(void* ptr, uint32_t value, } hsa_status_t BlitKernel::EnableProfiling(bool enable) { - core::Queue* cmd_queue = core::Queue::Convert(queue_); - if (cmd_queue != NULL) { - AMD_HSA_BITS_SET(cmd_queue->amd_queue_.queue_properties, - AMD_QUEUE_PROPERTIES_ENABLE_PROFILING, enable); - return HSA_STATUS_SUCCESS; - } - - return HSA_STATUS_ERROR_OUT_OF_RESOURCES; + AMD_HSA_BITS_SET(queue_->amd_queue_.queue_properties, + AMD_QUEUE_PROPERTIES_ENABLE_PROFILING, enable); + return HSA_STATUS_SUCCESS; } uint64_t BlitKernel::AcquireWriteIndex(uint32_t num_packet) { - assert(queue_->size >= num_packet); + assert(queue_->public_handle()->size >= num_packet); - uint64_t write_index = - HSA::hsa_queue_add_write_index_acq_rel(queue_, num_packet); + uint64_t write_index = queue_->AddWriteIndexAcqRel(num_packet); while (true) { // Wait until we have room in the queue; - const uint64_t read_index = HSA::hsa_queue_load_read_index_relaxed(queue_); - if ((write_index - read_index) < queue_->size) { + const uint64_t read_index = queue_->LoadReadIndexRelaxed(); + if ((write_index - read_index) <= queue_->public_handle()->size) { break; } } @@ -831,19 +799,10 @@ uint64_t BlitKernel::AcquireWriteIndex(uint32_t num_packet) { } void BlitKernel::ReleaseWriteIndex(uint64_t write_index, uint32_t num_packet) { - // Launch packet. - while (true) { - // Make sure that the address before ::current_offset is already released. - // Otherwise the packet processor may read invalid packets. - uint64_t expected_offset = write_index; - if (atomic::Cas(&cached_index_, write_index + num_packet, expected_offset, - std::memory_order_release) == expected_offset) { - // Update doorbel register with last packet id. - HSA::hsa_signal_store_release(queue_->doorbell_signal, - write_index + num_packet - 1); - break; - } - } + // Update doorbel register with last packet id. + core::Signal* doorbell = + core::Signal::Convert(queue_->public_handle()->doorbell_signal); + doorbell->StoreRelease(write_index + num_packet - 1); } hsa_status_t BlitKernel::FenceRelease(uint64_t write_index, @@ -871,7 +830,8 @@ hsa_status_t BlitKernel::FenceRelease(uint64_t write_index, // Populate queue buffer with AQL packet. hsa_barrier_and_packet_t* queue_buffer = - reinterpret_cast(queue_->base_address); + reinterpret_cast( + queue_->public_handle()->base_address); std::atomic_thread_fence(std::memory_order_acquire); queue_buffer[(write_index + num_copy_packet) & queue_bitmask_] = packet; std::atomic_thread_fence(std::memory_order_release); @@ -921,7 +881,8 @@ void BlitKernel::PopulateQueue(uint64_t index, uint64_t code_handle, void* args, // Populate queue buffer with AQL packet. hsa_kernel_dispatch_packet_t* queue_buffer = - reinterpret_cast(queue_->base_address); + reinterpret_cast( + queue_->public_handle()->base_address); std::atomic_thread_fence(std::memory_order_acquire); queue_buffer[index & queue_bitmask_] = packet; std::atomic_thread_fence(std::memory_order_release); diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp index d56f7fb518..20e4dc9000 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp @@ -71,9 +71,8 @@ GpuAgent::GpuAgent(HSAuint32 node, const HsaNodeProperties& node_props) : GpuAgentInt(node), properties_(node_props), current_coherency_type_(HSA_AMD_COHERENCY_TYPE_COHERENT), - blit_h2d_(NULL), - blit_d2h_(NULL), - blit_d2d_(NULL), + blits_(), + queues_(), local_region_(NULL), is_kv_device_(false), trap_code_buf_(NULL), @@ -135,28 +134,16 @@ GpuAgent::GpuAgent(HSAuint32 node, const HsaNodeProperties& node_props) } GpuAgent::~GpuAgent() { - if (blit_h2d_ != NULL) { - hsa_status_t status = blit_h2d_->Destroy(*this); - assert(status == HSA_STATUS_SUCCESS); - - delete blit_h2d_; - blit_h2d_ = NULL; + for (int i = 0; i < BlitCount; ++i) { + if (blits_[i] != NULL) { + hsa_status_t status = blits_[i]->Destroy(*this); + assert(status == HSA_STATUS_SUCCESS); + delete blits_[i]; + } } - if (blit_d2h_ != NULL && blit_d2h_ != blit_d2d_) { - hsa_status_t status = blit_d2h_->Destroy(*this); - assert(status == HSA_STATUS_SUCCESS); - - delete blit_d2h_; - blit_d2h_ = NULL; - } - - if (blit_d2d_ != NULL) { - hsa_status_t status = blit_d2d_->Destroy(*this); - assert(status == HSA_STATUS_SUCCESS); - - delete blit_d2d_; - blit_d2d_ = NULL; + for (int i = 0; i < QueueCount; ++i) { + delete queues_[i]; } if (end_ts_base_addr_ != NULL) { @@ -501,6 +488,21 @@ hsa_status_t GpuAgent::VisitRegion( return HSA_STATUS_SUCCESS; } +core::Queue* GpuAgent::CreateInterceptibleQueue() { + // Until tools runtime is merged in we need to use HSA API + // rather than GpuAgent::QueueCreate to allow interception. + hsa_queue_t* queue_handle; + hsa_status_t status = + HSA::hsa_queue_create(public_handle(), minAqlSize_, HSA_QUEUE_TYPE_MULTI, + NULL, NULL, 0, 0, &queue_handle); + + if (status != HSA_STATUS_SUCCESS) { + return NULL; + } + + return core::Queue::Convert(queue_handle); +} + core::Blit* GpuAgent::CreateBlitSdma() { BlitSdma* sdma = new BlitSdma(); @@ -513,8 +515,8 @@ core::Blit* GpuAgent::CreateBlitSdma() { return sdma; } -core::Blit* GpuAgent::CreateBlitKernel() { - BlitKernel* kernl = new BlitKernel(); +core::Blit* GpuAgent::CreateBlitKernel(core::Queue* queue) { + BlitKernel* kernl = new BlitKernel(queue); if (kernl->Initialize(*this) != HSA_STATUS_SUCCESS) { kernl->Destroy(*this); @@ -536,25 +538,29 @@ void GpuAgent::InitDma() { // Try create SDMA blit first. if (core::Runtime::runtime_singleton_->flag().enable_sdma() && (HSA_PROFILE_BASE == profile_)) { - blit_h2d_ = CreateBlitSdma(); - blit_d2h_ = CreateBlitSdma(); + blits_[BlitHostToDev] = CreateBlitSdma(); + blits_[BlitDevToHost] = CreateBlitSdma(); - if (blit_h2d_ != NULL && blit_d2h_ != NULL) { + if (blits_[BlitHostToDev] != NULL && blits_[BlitDevToHost] != NULL) { blit_initialized_.store(true, std::memory_order_release); return; } } // Fall back to blit kernel if SDMA is unavailable. - assert(blit_h2d_ == NULL || blit_d2h_ == NULL); + if (blits_[BlitHostToDev] == NULL) { + // Create a dedicated compute queue for host-to-device blits. + queues_[QueueBlitOnly] = CreateInterceptibleQueue(); + assert(queues_[QueueBlitOnly] != NULL && "Queue creation failed"); - if (blit_h2d_ == NULL) { - blit_h2d_ = CreateBlitKernel(); + blits_[BlitHostToDev] = CreateBlitKernel(queues_[QueueBlitOnly]); + assert(blits_[BlitHostToDev] != NULL && "Blit creation failed"); } - if (blit_d2h_ == NULL) { - // Share device-to-host queue with device-to-device. - blit_d2h_ = blit_d2d_; + if (blits_[BlitDevToHost] == NULL) { + // Share utility queue with device-to-host blits. + blits_[BlitDevToHost] = CreateBlitKernel(queues_[QueueUtility]); + assert(blits_[BlitDevToHost] != NULL && "Blit creation failed"); } blit_initialized_.store(true, std::memory_order_release); @@ -562,23 +568,26 @@ void GpuAgent::InitDma() { } } -hsa_status_t GpuAgent::InitBlitKernel() { - // Unlike InitDma, this function is not designed for lazy initialization. - // So checking the state without double checked locking is fine. - if (blit_d2d_ == NULL) { - blit_d2d_ = CreateBlitKernel(); - } +hsa_status_t GpuAgent::PostToolsInit() { + // Defer utility queue creation to allow tools to intercept. + queues_[QueueUtility] = CreateInterceptibleQueue(); - return (blit_d2d_ != NULL) ? HSA_STATUS_SUCCESS - : HSA_STATUS_ERROR_OUT_OF_RESOURCES; -} - -hsa_status_t GpuAgent::DmaCopy(void* dst, const void* src, size_t size) { - if (blit_d2d_ == NULL) { + if (queues_[QueueUtility] == NULL) { return HSA_STATUS_ERROR_OUT_OF_RESOURCES; } - return blit_d2d_->SubmitLinearCopyCommand(dst, src, size); + // Share utility queue with device-to-device blits. + blits_[BlitDevToDev] = CreateBlitKernel(queues_[QueueUtility]); + + if (blits_[BlitDevToDev] == NULL) { + return HSA_STATUS_ERROR_OUT_OF_RESOURCES; + } + + return HSA_STATUS_SUCCESS; +} + +hsa_status_t GpuAgent::DmaCopy(void* dst, const void* src, size_t size) { + return blits_[BlitDevToDev]->SubmitLinearCopyCommand(dst, src, size); } hsa_status_t GpuAgent::DmaCopy(void* dst, core::Agent& dst_agent, @@ -589,11 +598,11 @@ hsa_status_t GpuAgent::DmaCopy(void* dst, core::Agent& dst_agent, core::Blit* blit = (src_agent.device_type() == core::Agent::kAmdCpuDevice && dst_agent.device_type() == core::Agent::kAmdGpuDevice) - ? blit_h2d_ + ? blits_[BlitHostToDev] : (src_agent.device_type() == core::Agent::kAmdGpuDevice && dst_agent.device_type() == core::Agent::kAmdCpuDevice) - ? blit_d2h_ - : blit_d2d_; + ? blits_[BlitDevToHost] + : blits_[BlitDevToDev]; if (blit == NULL) { return HSA_STATUS_ERROR_OUT_OF_RESOURCES; @@ -612,11 +621,7 @@ hsa_status_t GpuAgent::DmaCopy(void* dst, core::Agent& dst_agent, } hsa_status_t GpuAgent::DmaFill(void* ptr, uint32_t value, size_t count) { - if (blit_d2d_ == NULL) { - return HSA_STATUS_ERROR_OUT_OF_RESOURCES; - } - - return blit_d2d_->SubmitLinearFillCommand(ptr, value, count); + return blits_[BlitDevToDev]->SubmitLinearFillCommand(ptr, value, count); } hsa_status_t GpuAgent::EnableDmaProfiling(bool enable) { @@ -624,10 +629,9 @@ hsa_status_t GpuAgent::EnableDmaProfiling(bool enable) { return HSA_STATUS_ERROR_OUT_OF_RESOURCES; } - core::Blit* blit[3] = {blit_h2d_, blit_d2h_, blit_d2d_}; - for (int i = 0; i < 3; ++i) { - if (blit[i] != NULL) { - const hsa_status_t stat = blit[i]->EnableProfiling(enable); + for (int i = 0; i < BlitCount; ++i) { + if (blits_[i] != NULL) { + const hsa_status_t stat = blits_[i]->EnableProfiling(enable); if (stat != HSA_STATUS_SUCCESS) { return stat; } diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp index a93d75c99e..26fb117b56 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp @@ -821,11 +821,9 @@ void Runtime::Load() { // Load tools libraries LoadTools(); - // Initialize blit kernel object after tools is initialized to allow tools - // to overload blit kernel. for (core::Agent* agent : gpu_agents_) { const hsa_status_t stat = - reinterpret_cast(agent)->InitBlitKernel(); + reinterpret_cast(agent)->PostToolsInit(); assert(HSA_STATUS_SUCCESS == stat); } }