lmoriche · Nov 5, 2026 · shwetagkhatri · Nov 5, 2026 · cfreeamd · Nov 5, 2026
@@ -1292,12 +1292,14 @@ bool Device::populateOCLDeviceConstants() {
                       static_cast<hsa_agent_info_t>(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_agent_info_t>(HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED),
-                        &signal_hints_);
+                        static_cast<hsa_agent_info_t>(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, "

@@ -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<const amd_queue_t*>(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_agent_info_t>(
-                                   HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED),
-                               &hints_supported),
+                               static_cast<hsa_agent_info_t>(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<uint32_t>(aql_extensions[0]);
   if (!hints_supported) {
     markAsSkip();
     std::cout << "Test not applicable as the GPU's CP does not support signal hints. Skipping."

@@ -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);

@@ -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_agent_info_t>(
-                                   HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED),
-                               &hints_supported),
+                               static_cast<hsa_agent_info_t>(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;

@@ -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);

@@ -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<bool*>(value) = HostAtomicsSupported();
       break;
-    case HSA_AMD_AGENT_INFO_SIGNAL_HINTS_SUPPORTED:
-      *static_cast<bool*>(value) = SignalHintsSupported();
-      break;
     default:
       return HSA_STATUS_ERROR_INVALID_ARGUMENT;
       break;

@@ -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 {

@@ -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