ROCm 1.6.1 updates.

Author: jedwards-AMD

src/core/inc/queue.h

@@ -64,15 +64,13 @@ struct AqlPacket {
     hsa_agent_dispatch_packet_t agent;
   };
 
-  uint8_t type() {
+  uint8_t type() const {
     return ((dispatch.header >> HSA_PACKET_HEADER_TYPE) &
                       ((1 << HSA_PACKET_HEADER_WIDTH_TYPE) - 1));
   }
 
-  bool IsValid() {
-    const uint8_t packet_type = dispatch.header >> HSA_PACKET_HEADER_TYPE;
-    return (packet_type > HSA_PACKET_TYPE_INVALID &&
-            packet_type <= HSA_PACKET_TYPE_BARRIER_OR);
+  bool IsValid() const {
+    return (type() <= HSA_PACKET_TYPE_BARRIER_OR) & (type() != HSA_PACKET_TYPE_INVALID);
   }
 
   std::string string() const {

src/core/runtime/amd_aql_queue.cpp

@@ -775,17 +775,20 @@ void AqlQueue::ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b) {
   // pm4_ib_buf_ is a shared resource, so mutually exclude here.
   ScopedAcquire<KernelMutex> lock(&pm4_ib_mutex_);
 
+  // Obtain reference to any container queue.
+  core::Queue* queue = core::Queue::Convert(public_handle());
+
   // Obtain a queue slot for a single AQL packet.
-  uint64_t write_idx = AddWriteIndexAcqRel(1);
+  uint64_t write_idx = queue->AddWriteIndexAcqRel(1);
 
-  while ((write_idx - LoadReadIndexRelaxed()) > public_handle()->size) {
+  while ((write_idx - queue->LoadReadIndexRelaxed()) > queue->amd_queue_.hsa_queue.size) {
     os::YieldThread();
   }
 
-  uint32_t slot_idx = uint32_t(write_idx % public_handle()->size);
+  uint32_t slot_idx = uint32_t(write_idx % queue->amd_queue_.hsa_queue.size);
   constexpr uint32_t slot_size_b = 0x40;
   uint32_t* queue_slot =
-      (uint32_t*)(uintptr_t(public_handle()->base_address) + (slot_idx * slot_size_b));
+      (uint32_t*)(uintptr_t(queue->amd_queue_.hsa_queue.base_address) + (slot_idx * slot_size_b));
 
   // Copy client PM4 command into IB.
   assert(cmd_size_b < pm4_ib_size_b_ && "PM4 exceeds IB size");
@@ -877,12 +880,13 @@ void AqlQueue::ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b) {
   atomic::Store(&queue_slot[0], slot_data[0], std::memory_order_release);
 
   // Submit the packet slot.
-  core::Signal* doorbell =
-      core::Signal::Convert(public_handle()->doorbell_signal);
+  core::Signal* doorbell = core::Signal::Convert(queue->amd_queue_.hsa_queue.doorbell_signal);
   doorbell->StoreRelease(write_idx);
 
   // Wait for the packet to be consumed.
-  while (LoadReadIndexRelaxed() <= write_idx) {
+  // Should be switched to a signal wait when aql_pm4_ib can be used on all
+  // supported platforms.
+  while (queue->LoadReadIndexRelaxed() <= write_idx) {
     os::YieldThread();
   }
 }

src/core/runtime/amd_gpu_agent.cpp

@@ -512,18 +512,10 @@ hsa_status_t GpuAgent::VisitRegion(
 }
 
 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);
+  // Disabled intercept of internal queues pending tools updates.
+  core::Queue* queue = nullptr;
+  QueueCreate(minAqlSize_, HSA_QUEUE_TYPE_MULTI, NULL, NULL, 0, 0, &queue);
+  return queue;
 }
 
 core::Blit* GpuAgent::CreateBlitSdma() {
@@ -643,15 +635,14 @@ hsa_status_t GpuAgent::DmaCopy(void* dst, core::Agent& dst_agent,
     return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
   }
 
-  hsa_status_t stat =
-      blit->SubmitLinearCopyCommand(dst, src, size, dep_signals, out_signal);
-
-  if (profiling_enabled() && HSA_STATUS_SUCCESS == stat) {
+  if (profiling_enabled()) {
     // Track the agent so we could translate the resulting timestamp to system
     // domain correctly.
     out_signal.async_copy_agent(this);
   }
 
+  hsa_status_t stat = blit->SubmitLinearCopyCommand(dst, src, size, dep_signals, out_signal);
+
   return stat;
 }
 

src/core/runtime/runtime.cpp

@@ -113,6 +113,7 @@ hsa_status_t Runtime::Acquire() {
     hsa_status_t status = runtime_singleton_->Load();
 
     if (status != HSA_STATUS_SUCCESS) {
+      runtime_singleton_->ref_count_--;
       return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
     }
   }
@@ -683,7 +684,7 @@ hsa_status_t Runtime::PtrInfo(void* ptr, hsa_amd_pointer_info_t* info, void* (*a
   hsa_amd_pointer_info_t retInfo;
 
   // check output struct is at least as large as the first info revision.
-  if (info->size < sizeof(struct hsa_amd_pointer_info_v1_s)) return HSA_STATUS_ERROR_INVALID_ARGUMENT;
+  if (info->size < 40) return HSA_STATUS_ERROR_INVALID_ARGUMENT;
 
   bool returnListData =
       ((alloc != nullptr) && (num_agents_accessible != nullptr) && (accessible != nullptr));

src/core/util/lnx/os_linux.cpp

@@ -322,11 +322,14 @@ int ResetOsEvent(EventHandle event) {
   return ret_code;
 }
 
+static double invPeriod = 0.0;
+
 uint64_t ReadAccurateClock() {
+  if (invPeriod == 0.0) AccurateClockFrequency();
   timespec time;
   int err = clock_gettime(CLOCK_MONOTONIC_RAW, &time);
   assert(err == 0 && "clock_gettime(CLOCK_MONOTONIC_RAW,...) failed");
-  return uint64_t(time.tv_sec) * 1000000000ull + uint64_t(time.tv_nsec);
+  return (uint64_t(time.tv_sec) * 1000000000ull + uint64_t(time.tv_nsec)) * invPeriod;
 }
 
 uint64_t AccurateClockFrequency() {
@@ -339,7 +342,8 @@ uint64_t AccurateClockFrequency() {
   assert(time.tv_nsec < 0xFFFFFFFF &&
          "clock_getres(CLOCK_MONOTONIC_RAW,...) returned very low frequency "
          "(<1Hz).");
-  return uint64_t(time.tv_nsec) * 1000000000ull;
+  if (invPeriod == 0.0) invPeriod = 1.0 / double(time.tv_nsec);
+  return 1000000000ull / uint64_t(time.tv_nsec);
 }
 }
 

src/core/util/timer.cpp

@@ -93,6 +93,8 @@ fast_clock::init::init() {
 
   fast_clock::freq = double(min) / duration_in_seconds(elapsed);
   fast_clock::period_ps = 1e12 / fast_clock::freq;
+  // printf("Timer setup took %f ms\n", duration_in_seconds(elapsed)*1000.0f);
+  // printf("Fast clock frequency: %f MHz\n", double(fast_clock::freq)/1e6);
 }
 
 double accurate_clock::period_ns;

src/inc/hsa_ext_amd.h

@@ -1320,7 +1320,7 @@ typedef enum {
  * @brief Describes a memory allocation known to ROCr.
  * Within a ROCr major version this structure can only grow.
  */
-typedef struct hsa_amd_pointer_info_v1_s {
+typedef struct hsa_amd_pointer_info_s {
   /*
   Size in bytes of this structure.  Used for version control within a major ROCr
   revision.  Set to sizeof(hsa_amd_pointer_t) prior to calling
@@ -1349,17 +1349,6 @@ typedef struct hsa_amd_pointer_info_v1_s {
   Application provided value.
   */
   void* userData;
-} hsa_amd_pointer_info_v1_t;
-
-/**
- * @brief Minor version updates to pointer info.
- */
-#ifdef __cplusplus
-typedef struct hsa_amd_pointer_info_v2_s : hsa_amd_pointer_info_v1_t {
-#else
-typedef struct hsa_amd_pointer_info_v2_t {
-  struct hsa_amd_pointer_info_v1_t;
-#endif
   /*
   Reports an agent which "owns" (ie has preferred access to) the pool in which the allocation was
   made.  When multiple agents share equal access to a pool (ex: multiple CPU agents, or multi-die