From 721e56ef5ca88103ec7cc894db4b71306321879b Mon Sep 17 00:00:00 2001 From: David Yat Sin Date: Wed, 30 Aug 2023 17:31:13 +0000 Subject: [PATCH] Extend ExecutePM4() to accept completion signal and fences ExecutePM4() function can optionally accept extra arguments for acquire fence scope, release fence scope andcompletion signal. When a completion signal is provided, ExecutePM4() does not wait for the commands to complete. Change-Id: Ib2a433b7bce1cb6260be8b76fe902335bd5dfada --- runtime/hsa-runtime/core/inc/amd_aql_queue.h | 7 +++- runtime/hsa-runtime/core/inc/host_queue.h | 5 ++- .../hsa-runtime/core/inc/intercept_queue.h | 7 +++- runtime/hsa-runtime/core/inc/queue.h | 20 ++++++++- .../core/runtime/amd_aql_queue.cpp | 42 ++++++++++++------- 5 files changed, 59 insertions(+), 22 deletions(-) diff --git a/runtime/hsa-runtime/core/inc/amd_aql_queue.h b/runtime/hsa-runtime/core/inc/amd_aql_queue.h index ddb8671c2d..7278d0b6ce 100644 --- a/runtime/hsa-runtime/core/inc/amd_aql_queue.h +++ b/runtime/hsa-runtime/core/inc/amd_aql_queue.h @@ -196,8 +196,11 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Doo /// @return hsa_status_t hsa_status_t GetCUMasking(uint32_t num_cu_mask_count, uint32_t* cu_mask) override; - /// @brief Submits a block of PM4 and waits until it has been executed. - void ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b) override; + // @brief Submits a block of PM4 and waits until it has been executed. + void ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b, + hsa_fence_scope_t acquireFence = HSA_FENCE_SCOPE_NONE, + hsa_fence_scope_t releaseFence = HSA_FENCE_SCOPE_NONE, + hsa_signal_t* signal = NULL) override; /// @brief Enables/Disables profiling overrides SetProfiling from core::Queue void SetProfiling(bool enabled) override; diff --git a/runtime/hsa-runtime/core/inc/host_queue.h b/runtime/hsa-runtime/core/inc/host_queue.h index 8521aed7b7..e16d277f3a 100644 --- a/runtime/hsa-runtime/core/inc/host_queue.h +++ b/runtime/hsa-runtime/core/inc/host_queue.h @@ -152,7 +152,10 @@ class HostQueue : public Queue { return HSA_STATUS_ERROR_INVALID_QUEUE; } - void ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b) override { + void ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b, + hsa_fence_scope_t acquireFence = HSA_FENCE_SCOPE_NONE, + hsa_fence_scope_t releaseFence = HSA_FENCE_SCOPE_NONE, + hsa_signal_t* signal = NULL) override { assert(false && "HostQueue::ExecutePM4 is unimplemented"); } diff --git a/runtime/hsa-runtime/core/inc/intercept_queue.h b/runtime/hsa-runtime/core/inc/intercept_queue.h index e7784ffaa5..f6bb483170 100644 --- a/runtime/hsa-runtime/core/inc/intercept_queue.h +++ b/runtime/hsa-runtime/core/inc/intercept_queue.h @@ -120,8 +120,11 @@ class QueueWrapper : public Queue { hsa_status_t GetCUMasking(uint32_t num_cu_mask_count, uint32_t* cu_mask) override { return wrapped->GetCUMasking(num_cu_mask_count, cu_mask); } - void ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b) override { - wrapped->ExecutePM4(cmd_data, cmd_size_b); + void ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b, + hsa_fence_scope_t acquireFence = HSA_FENCE_SCOPE_NONE, + hsa_fence_scope_t releaseFence = HSA_FENCE_SCOPE_NONE, + hsa_signal_t* signal = NULL) override { + wrapped->ExecutePM4(cmd_data, cmd_size_b, acquireFence, releaseFence, signal); } void SetProfiling(bool enabled) override { wrapped->SetProfiling(enabled); } diff --git a/runtime/hsa-runtime/core/inc/queue.h b/runtime/hsa-runtime/core/inc/queue.h index 1125102344..c3b5f0a279 100644 --- a/runtime/hsa-runtime/core/inc/queue.h +++ b/runtime/hsa-runtime/core/inc/queue.h @@ -346,8 +346,24 @@ class Queue : public Checked<0xFA3906A679F9DB49>, private LocalQueue { /// @return hsa_status_t virtual hsa_status_t GetCUMasking(uint32_t num_cu_mask_count, uint32_t* cu_mask) = 0; - // @brief Submits a block of PM4 and waits until it has been executed. - virtual void ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b) = 0; + /// @brief Submits a block of PM4. + /// + /// @param cmd_data pointer to command buffer + /// + /// @param cmd_size_b command buffer size in bytes + /// + /// @param acquireFence acquire-fence type + /// + /// @param releaseFence acquire-fence type + /// + /// @param signal optional wait signal + /// + /// if @p signal is provided, function will return without waiting for commands to be executed + /// if @p signal is NULL, waits until commands have been executed. + virtual void ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b, + hsa_fence_scope_t acquireFence = HSA_FENCE_SCOPE_NONE, + hsa_fence_scope_t releaseFence = HSA_FENCE_SCOPE_NONE, + hsa_signal_t* signal = NULL) = 0; virtual void SetProfiling(bool enabled) { AMD_HSA_BITS_SET(amd_queue_.queue_properties, AMD_QUEUE_PROPERTIES_ENABLE_PROFILING, diff --git a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp index 83ed5cdd26..79b6d72b34 100644 --- a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp @@ -1395,7 +1395,12 @@ void AqlQueue::SetProfiling(bool enabled) { return; } -void AqlQueue::ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b) { +// If in_signal is NULL then this ExecutePM4 will block and wait for PM4 commands to complete +// If in_signal is provided, then ExecutePM4 will return and caller may wait for in_signal +// Note: On gfx8, there is no completion signal support, so ExecutePM4 will block even if +// in_signal is provided, and it is still valid to check in_signal after ExecutePM4 returns. +void AqlQueue::ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b, hsa_fence_scope_t acquireFence, + hsa_fence_scope_t releaseFence, hsa_signal_t* in_signal) { // pm4_ib_buf_ is a shared resource, so mutually exclude here. ScopedAcquire lock(&pm4_ib_mutex_); @@ -1431,7 +1436,7 @@ void AqlQueue::ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b) { // To respect multi-producer semantics, first buffer commands for the queue slot. constexpr uint32_t slot_size_dw = uint32_t(slot_size_b / sizeof(uint32_t)); uint32_t slot_data[slot_size_dw]; - hsa_signal_t signal = {0}; + hsa_signal_t local_signal = {0}; hsa_status_t err; if (agent_->isa()->GetMajorVersion() <= 8) { @@ -1476,28 +1481,32 @@ void AqlQueue::ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b) { // Construct an AQL packet to jump to the PM4 IB. struct amd_aql_pm4_ib { uint16_t header; - uint8_t amd_format; - uint8_t reserved0; + uint16_t ven_hdr; uint32_t ib_jump_cmd[4]; uint32_t dw_cnt_remain; - uint32_t reserved1[8]; + uint32_t reserved[8]; hsa_signal_t completion_signal; }; + if (!in_signal) { + err = hsa_signal_create(1, 0, NULL, &local_signal); + assert(err == HSA_STATUS_SUCCESS); + } + constexpr uint32_t AMD_AQL_FORMAT_PM4_IB = 0x1; - err = hsa_signal_create(1, 0, NULL, &signal); - assert(err == HSA_STATUS_SUCCESS); - amd_aql_pm4_ib aql_pm4_ib{}; - aql_pm4_ib.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; - aql_pm4_ib.amd_format = AMD_AQL_FORMAT_PM4_IB; + aql_pm4_ib.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE | + (acquireFence << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) | + (releaseFence << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE); + + aql_pm4_ib.ven_hdr = AMD_AQL_FORMAT_PM4_IB; aql_pm4_ib.ib_jump_cmd[0] = ib_jump_cmd[0]; aql_pm4_ib.ib_jump_cmd[1] = ib_jump_cmd[1]; aql_pm4_ib.ib_jump_cmd[2] = ib_jump_cmd[2]; aql_pm4_ib.ib_jump_cmd[3] = ib_jump_cmd[3]; aql_pm4_ib.dw_cnt_remain = 0xA; - aql_pm4_ib.completion_signal = signal; + aql_pm4_ib.completion_signal = in_signal ? *in_signal : local_signal; memcpy(slot_data, &aql_pm4_ib, sizeof(aql_pm4_ib)); } else { @@ -1518,11 +1527,14 @@ void AqlQueue::ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b) { if (agent_->isa()->GetMajorVersion() <= 8) { while (queue->LoadReadIndexRelaxed() <= write_idx) os::YieldThread(); - } else { + + if (in_signal) hsa_signal_store_screlease(*in_signal, 0); + } else if (!in_signal) { + // On gfx9 and newer, if in_signal is not provided, we block and wait for own signal hsa_signal_value_t ret; - ret = hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, 1, - (uint64_t)-1, HSA_WAIT_STATE_ACTIVE); - err = hsa_signal_destroy(signal); + ret = hsa_signal_wait_scacquire(local_signal, HSA_SIGNAL_CONDITION_LT, 1, (uint64_t)-1, + HSA_WAIT_STATE_ACTIVE); + err = hsa_signal_destroy(local_signal); assert(ret == 0 && err == HSA_STATUS_SUCCESS); } }