diff --git a/projects/clr/rocclr/device/rocm/rocdevice.cpp b/projects/clr/rocclr/device/rocm/rocdevice.cpp index bc8e5a7178f5..9a3c1c99b635 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.cpp @@ -1292,12 +1292,14 @@ bool Device::populateOCLDeviceConstants() { static_cast(HSA_AMD_AGENT_INFO_PCIE_STORES_NEED_FLUSH), &pcie_stores_need_flush_); // CP firmware without hint support takes a hinted handle as the signal address. Left false by - // a runtime without the query. A profiler tool hands packet completion signals to ROCr, which - // rejects a hinted handle. + // a runtime that does not report the AQL extension. A profiler tool hands packet completion + // signals to ROCr, which rejects a hinted handle. if (!tool_attached) { + uint8_t aql_extensions[8] = {}; Hsa::agent_get_info(bkendDevice_, - static_cast(HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED), - &signal_hints_); + static_cast(HSA_AMD_AGENT_INFO_AQL_EXTENSIONS), + aql_extensions); + signal_hints_ = hsa_flag_isset64(aql_extensions, HSA_AMD_AQL_PROPERTY_SIGNAL_HINTS); } ClPrint(amd::LOG_INFO, amd::LOG_INIT, "Signal pool: %s. Device memory: graph signals %d, queue signals %d, " diff --git a/projects/rocr-runtime/rocrtst/suites/functional/signal_hints.cc b/projects/rocr-runtime/rocrtst/suites/functional/signal_hints.cc index 2765d96fe341..7760caa157a7 100644 --- a/projects/rocr-runtime/rocrtst/suites/functional/signal_hints.cc +++ b/projects/rocr-runtime/rocrtst/suites/functional/signal_hints.cc @@ -15,6 +15,7 @@ #include "common/base_rocr_utils.h" #include "common/common.h" #include "gtest/gtest.h" +#include "hsa/amd_hsa_queue.h" #include "hsa/amd_hsa_signal.h" #include "hsa/hsa_ext_amd.h" @@ -44,6 +45,24 @@ hsa_signal_t CreateHostSignal(hsa_signal_value_t initial_value) { return signal; } +// Whether the CP reports AMD_QUEUE_CAPS_CP_SIGNAL_HINTS in a queue's caps, which it writes once it +// has processed a packet on the queue +bool CpReportsSignalHints(hsa_agent_t agent) { + hsa_queue_t* queue = nullptr; + EXPECT_EQ(rocrtst::CreateQueue(agent, &queue, kQueueSize), HSA_STATUS_SUCCESS); + if (queue == nullptr) return false; + const hsa_signal_t completion = CreateHostSignal(1); + RingDoorbell(queue, rocrtst::WriteBarrierAndToQueue(queue, hsa_signal_t{0}, completion)); + EXPECT_EQ(hsa_signal_wait_scacquire(completion, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, + HSA_WAIT_STATE_ACTIVE), + 0); + // The queue is the hsa_queue_t at the start of the runtime's amd_queue_t + const uint32_t caps = reinterpret_cast(queue)->caps; + EXPECT_EQ(hsa_signal_destroy(completion), HSA_STATUS_SUCCESS); + EXPECT_EQ(hsa_queue_destroy(queue), HSA_STATUS_SUCCESS); + return (caps & AMD_QUEUE_CAPS_CP_SIGNAL_HINTS) != 0; +} + hsa_signal_t CreateDeviceSignal(hsa_amd_memory_pool_t pool, hsa_signal_value_t initial_value) { hsa_signal_t signal = {0}; hsa_amd_signal_batch_t batch; @@ -73,12 +92,15 @@ void SignalHintsTest::SetUp(void) { if (test_skipped_) return; ASSERT_EQ(rocrtst::SetDefaultAgents(this), HSA_STATUS_SUCCESS); - bool hints_supported = false; + uint8_t aql_extensions[8] = {}; ASSERT_EQ(hsa_agent_get_info(*gpu_device1(), - static_cast( - HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED), - &hints_supported), + static_cast(HSA_AMD_AGENT_INFO_AQL_EXTENSIONS), + aql_extensions), HSA_STATUS_SUCCESS); + const bool hints_supported = hsa_flag_isset64(aql_extensions, HSA_AMD_AQL_PROPERTY_SIGNAL_HINTS); + // ROCr sets the bit from the same caps, at the index hsa_flag_isset64 reads + EXPECT_EQ(hints_supported, CpReportsSignalHints(*gpu_device1())) + << "AQL_EXTENSIONS bytes " << std::hex << static_cast(aql_extensions[0]); if (!hints_supported) { markAsSkip(); std::cout << "Test not applicable as the GPU's CP does not support signal hints. Skipping." diff --git a/projects/rocr-runtime/rocrtst/suites/functional/signal_hints.h b/projects/rocr-runtime/rocrtst/suites/functional/signal_hints.h index 43580ecedfd8..8810b1d878a7 100644 --- a/projects/rocr-runtime/rocrtst/suites/functional/signal_hints.h +++ b/projects/rocr-runtime/rocrtst/suites/functional/signal_hints.h @@ -16,7 +16,7 @@ // Checks that the CP honors each completion signal handle hint (amd_hsa_signal.h) in the AQL // packets it processes. Skipped on an agent that does not report -// HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED. +// HSA_AMD_AQL_PROPERTY_SIGNAL_HINTS in HSA_AMD_AGENT_INFO_AQL_EXTENSIONS. class SignalHintsTest : public TestBase { public: SignalHintsTest(void); diff --git a/projects/rocr-runtime/rocrtst/suites/performance/signal_hint_cost.cc b/projects/rocr-runtime/rocrtst/suites/performance/signal_hint_cost.cc index b6f8298406aa..398424d8108a 100644 --- a/projects/rocr-runtime/rocrtst/suites/performance/signal_hint_cost.cc +++ b/projects/rocr-runtime/rocrtst/suites/performance/signal_hint_cost.cc @@ -65,13 +65,12 @@ void SignalHintCost::SetUp(void) { if (test_skipped_) return; ASSERT_EQ(rocrtst::SetDefaultAgents(this), HSA_STATUS_SUCCESS); - bool hints_supported = false; + uint8_t aql_extensions[8] = {}; ASSERT_EQ(hsa_agent_get_info(*gpu_device1(), - static_cast( - HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED), - &hints_supported), + static_cast(HSA_AMD_AGENT_INFO_AQL_EXTENSIONS), + aql_extensions), HSA_STATUS_SUCCESS); - if (!hints_supported) { + if (!hsa_flag_isset64(aql_extensions, HSA_AMD_AQL_PROPERTY_SIGNAL_HINTS)) { markAsSkip(); std::cout << "Test not applicable as the GPU's CP does not support signal hints. Skipping." << std::endl; diff --git a/projects/rocr-runtime/rocrtst/suites/performance/signal_hint_cost.h b/projects/rocr-runtime/rocrtst/suites/performance/signal_hint_cost.h index d83fb19e2222..646d4debb0a6 100644 --- a/projects/rocr-runtime/rocrtst/suites/performance/signal_hint_cost.h +++ b/projects/rocr-runtime/rocrtst/suites/performance/signal_hint_cost.h @@ -18,7 +18,7 @@ // Measures the CP time per barrier packet that completes a signal, with each completion signal // handle hint and with all of them, for signals in host memory, device memory and with // interrupt events. Skipped on an agent that does not report -// HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED. +// HSA_AMD_AQL_PROPERTY_SIGNAL_HINTS in HSA_AMD_AGENT_INFO_AQL_EXTENSIONS. class SignalHintCost : public TestBase { public: SignalHintCost(void); 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 7762051fb26d..c3143882f230 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 @@ -2368,6 +2368,10 @@ void GpuAgent::GetAqlInfoProperties(uint8_t value[8]) const { // Fill the HSA_AMD_AQL_PROPERTY_EXT_DISPATCH if (extended_aql_dispatch_supported_) setFlag(HSA_AMD_AQL_PROPERTY_EXT_DISPATCH); + + // Fill the HSA_AMD_AQL_PROPERTY_SIGNAL_HINTS + if (SignalHintsSupported()) + setFlag(HSA_AMD_AQL_PROPERTY_SIGNAL_HINTS); } @@ -2805,9 +2809,6 @@ hsa_status_t GpuAgent::GetInfo(hsa_agent_info_t attribute, void* value) const { case HSA_AMD_AGENT_INFO_HOST_ATOMICS_SUPPORTED: *static_cast(value) = HostAtomicsSupported(); break; - case HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED: - *static_cast(value) = SignalHintsSupported(); - break; default: return HSA_STATUS_ERROR_INVALID_ARGUMENT; break; diff --git a/projects/rocr-runtime/runtime/hsa-runtime/inc/amd_hsa_signal.h b/projects/rocr-runtime/runtime/hsa-runtime/inc/amd_hsa_signal.h index e6e45ceb356e..92f64bf02f9f 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/inc/amd_hsa_signal.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/inc/amd_hsa_signal.h @@ -81,8 +81,8 @@ typedef struct AMD_SIGNAL_ALIGN amd_signal_s { // completion signal without loading it. Hints live only in the handle, never // in the signal, so they may differ per use. Zero is today's behaviour. // Hints are gated on a CP capability and binding once advertised: a producer -// may set them only on an agent that reports -// HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED, and its CP firmware consumes every +// may set them only on an agent that reports HSA_AMD_AQL_PROPERTY_SIGNAL_HINTS +// in HSA_AMD_AGENT_INFO_AQL_EXTENSIONS, and its CP firmware consumes every // defined hint. Firmware without support takes the handle as the signal // address. Reserved hint bits must be zero. typedef enum { diff --git a/projects/rocr-runtime/runtime/hsa-runtime/inc/hsa_ext_amd.h b/projects/rocr-runtime/runtime/hsa-runtime/inc/hsa_ext_amd.h index d259714a102f..ca94f1a85f0b 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/inc/hsa_ext_amd.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/inc/hsa_ext_amd.h @@ -91,13 +91,17 @@ * - 1.36 - hsa_amd_signal_batch_create, hsa_amd_signal_batch_destroy, * HSA_AMD_MEMORY_POOL_INFO_SIGNAL_CAPABLE, * HSA_AMD_AGENT_INFO_PCIE_STORES_NEED_FLUSH, HSA_AMD_AGENT_INFO_HOST_ATOMICS_SUPPORTED, - * HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED, + * HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED (removed in 1.39, never released), * hsa_amd_signal_set_hints, hsa_amd_signal_get_hints (inline) * - 1.37 - hsa_amd_memory_async_batch_copy: HSA_AMD_MEMORY_COPY_OP_RECT * - 1.38 - hsa_amd_pointer_info: HSA_EXT_POINTER_TYPE_SVM + * - 1.39 - hsa_amd_agent_aql_properties_t: HSA_AMD_AQL_PROPERTY_SIGNAL_HINTS, which replaces + * HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED. 1.36-1.38 were never in a release + * (development and BKC snapshots only), so the 1.36 query is removed without a + * deprecation period. */ #define HSA_AMD_INTERFACE_VERSION_MAJOR 1 -#define HSA_AMD_INTERFACE_VERSION_MINOR 38 +#define HSA_AMD_INTERFACE_VERSION_MINOR 39 #ifdef __cplusplus extern "C" { @@ -911,7 +915,9 @@ typedef enum hsa_amd_agent_info_s { /** * Bit-mask indicating AQL Extensions supported by this agent. An AQL extension is set if the flag * bit is set at that position. User may use the hsa_flag_isset64 macro to verify whether a flag - * is set. The type of this attribute is uint8_t[8]. + * is set. The position of an extension is the value of its ::hsa_amd_agent_aql_properties_t + * enumerator, a bit index, not a mask. + * The type of this attribute is uint8_t[8]. */ HSA_AMD_AGENT_INFO_AQL_EXTENSIONS = 0xA115, /** @@ -1033,13 +1039,8 @@ typedef enum hsa_amd_agent_info_s { * The type of this attribute is bool. */ HSA_AMD_AGENT_INFO_HOST_ATOMICS_SUPPORTED = 0xA128, - /** - * Whether the agent's CP firmware consumes completion signal handle hints - * (::hsa_amd_signal_set_hints). Hints may be put in the completion_signal - * of an AQL packet only on an agent that reports true. GPU agents only. - * The type of this attribute is bool. - */ - HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED = 0xA129, + /* 0xA129 is not reused: it was HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED, + replaced by ::HSA_AMD_AQL_PROPERTY_SIGNAL_HINTS. */ } hsa_amd_agent_info_t; /** @@ -1051,9 +1052,22 @@ typedef enum hsa_amd_agent_memory_properties_s { /** * @brief Agent AQL info properties attributes + * + * Each enumerator's value is the index of its bit in + * ::HSA_AMD_AGENT_INFO_AQL_EXTENSIONS, not a mask: test it with + * hsa_flag_isset64(value, HSA_AMD_AQL_PROPERTY_*), which is how the runtime + * sets it (HSA_AMD_AQL_PROPERTY_SIGNAL_HINTS, 2, is bit 2 of the array). + * value[0] & HSA_AMD_AQL_PROPERTY_* tests another bit. */ typedef enum hsa_amd_agent_aql_properties_s { HSA_AMD_AQL_PROPERTY_EXT_DISPATCH = (1 << 0), + /** + * The agent's CP firmware consumes completion signal handle hints + * (::hsa_amd_signal_set_hints). Hints may be put in the completion_signal + * of an AQL packet only on an agent that reports this extension. GPU agents + * only. + */ + HSA_AMD_AQL_PROPERTY_SIGNAL_HINTS = (1 << 1), } hsa_amd_agent_aql_properties_t; /** @@ -2240,8 +2254,8 @@ hsa_status_t HSA_API hsa_amd_signal_batch_destroy(hsa_amd_signal_batch_t batch); * @brief Return @p signal with its handle hints replaced by @p hints. * * @p hints combines ::amd_signal_hint_t values; 0 returns the plain handle. - * Use hints only on an agent that reports - * ::HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED: its CP firmware consumes every + * Use hints only on an agent that reports ::HSA_AMD_AQL_PROPERTY_SIGNAL_HINTS + * in ::HSA_AMD_AGENT_INFO_AQL_EXTENSIONS: its CP firmware consumes every * defined hint, while firmware without support takes the hinted handle as the * signal address. Bits outside * ::AMD_SIGNAL_HINT_MASK are dropped. A hinted handle is only for the