Allow reducing max occupancy (max scratch waves) when applications request large amounts of scratch.

Also emit error messages to stderr if no async queue error callback was registered and queue fault messages are enabled (on by default). Queue fault messages are controlled with env key HSA_ENABLE_QUEUE_FAULT_MESSAGE. Change-Id: I496487b8d048b83aa95b9784e92928211f167b17 [ROCm/ROCR-Runtime commit: 0e17cc2887]
2016-12-20 05:50:15 -06:00
@@ -149,6 +149,7 @@ set ( SRCS "core/util/lnx/os_linux.cpp"
            "core/runtime/isa.cpp"
            "core/runtime/runtime.cpp"
            "core/runtime/signal.cpp"
+            "core/runtime/queue.cpp"
            "core/runtime/cache.cpp"
            "core/common/shared.cpp"
            "core/common/hsa_table_interface.cpp"
@@ -300,6 +300,9 @@ class Queue : public Checked<0xFA3906A679F9DB49>,
  // @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 Reports async queue errors to stderr if no other error handler was registered.
+  static void DefaultErrorHandler(hsa_status_t status, hsa_queue_t* source, void* data);
+
  // Handle of AMD Queue struct
  amd_queue_t& amd_queue_;

@@ -683,10 +683,7 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) {
    // Align whole waves to 1KB.
    scratch.size_per_thread = AlignUp(scratch.size_per_thread, 16);
    scratch.size = scratch.size_per_thread * (queue->amd_queue_.max_cu_id + 1) *
-                   32 * 64;  // TODO: replace constants.
-
-    // printf("Growing scratch to %u - %u\n", uint32_t(scratch.size_per_thread),
-    // uint32_t(scratch.size));
+        queue->agent_->properties().MaxSlotsScratchCU * queue->agent_->properties().WaveFrontSize;

    queue->agent_->AcquireQueueScratch(scratch);
    if (scratch.queue_base == NULL) {
@@ -921,18 +921,57 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) {

  ScopedAcquire<KernelMutex> lock(&scratch_lock_);
  scratch.queue_base = scratch_pool_.alloc(scratch.size);
-  scratch.queue_process_offset =
-      uintptr_t(scratch.queue_base) - uintptr_t(scratch_pool_.base());
+  scratch.queue_process_offset = uintptr_t(scratch.queue_base) - uintptr_t(scratch_pool_.base());

-  if ((scratch.queue_base != NULL) && (profile_ == HSA_PROFILE_BASE)) {
-    HSAuint64 alternate_va;
-    if (HSAKMT_STATUS_SUCCESS !=
-        hsaKmtMapMemoryToGPU(scratch.queue_base, scratch.size, &alternate_va)) {
-      assert(false && "Map scratch subrange failed!");
-      scratch_pool_.free(scratch.queue_base);
-      scratch.queue_base = NULL;
+  if (scratch.queue_base != NULL) {
+    if (profile_ == HSA_PROFILE_FULL) return;
+    if (profile_ == HSA_PROFILE_BASE) {
+      HSAuint64 alternate_va;
+      if (HSAKMT_STATUS_SUCCESS ==
+          hsaKmtMapMemoryToGPU(scratch.queue_base, scratch.size, &alternate_va))
+        return;
    }
  }
+
+  // Scratch request failed allocation or mapping.
+  scratch_pool_.free(scratch.queue_base);
+  scratch.queue_base = NULL;
+
+// Attempt to trim the maximum number of concurrent waves to allow scratch to fit.
+// This is somewhat dangerous as it limits the number of concurrent waves from future dispatches
+// on the queue if those waves use even small amounts of scratch.
+#ifndef NDEBUG
+  if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message())
+    fprintf(stderr, "Failed to map requested scratch - reducing queue occupancy.\n");
+#endif
+  uint64_t num_cus = properties_.NumFComputeCores / properties_.NumSIMDPerCU;
+  uint64_t size_per_wave = AlignUp(scratch.size_per_thread * properties_.WaveFrontSize, 1024);
+  uint64_t total_waves = scratch.size / size_per_wave;
+  uint64_t waves_per_cu = total_waves / num_cus;
+  while (waves_per_cu != 0) {
+    size_t size = waves_per_cu * num_cus * size_per_wave;
+    void* base = scratch_pool_.alloc(size);
+    HSAuint64 alternate_va;
+    if ((base != NULL) &&
+        ((profile_ == HSA_PROFILE_FULL) ||
+         (hsaKmtMapMemoryToGPU(base, size, &alternate_va) == HSAKMT_STATUS_SUCCESS))) {
+      // Scratch allocated and either full profile or map succeeded.
+      scratch.queue_base = base;
+      scratch.size = size;
+      scratch.queue_process_offset =
+          uintptr_t(scratch.queue_base) - uintptr_t(scratch_pool_.base());
+      return;
+    }
+    scratch_pool_.free(base);
+    waves_per_cu--;
+  }
+
+  // Failed to allocate minimal scratch
+  assert(scratch.queue_base == NULL && "bad scratch data");
+#ifndef NDEBUG
+  if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message())
+    fprintf(stderr, "Could not allocate scratch for one wave per CU.\n");
+#endif
 }

 void GpuAgent::ReleaseQueueScratch(void* base) {
@@ -281,6 +281,7 @@ hsa_status_t MemoryRegion::GetInfo(hsa_region_info_t attribute,
        case HSA_HEAPTYPE_FRAME_BUFFER_PRIVATE:
        case HSA_HEAPTYPE_FRAME_BUFFER_PUBLIC:
        case HSA_HEAPTYPE_SYSTEM:
+        case HSA_HEAPTYPE_GPU_SCRATCH:
          *((size_t*)value) = max_single_alloc_size_;
          break;
        default:
@@ -564,6 +564,8 @@ hsa_status_t hsa_queue_create(
    return HSA_STATUS_ERROR_INVALID_QUEUE_CREATION;
  }

+  if (callback == NULL) callback = core::Queue::DefaultErrorHandler;
+
  core::Queue* cmd_queue = NULL;
  status = agent->QueueCreate(size, type, callback, data, private_segment_size,
                              group_segment_size, &cmd_queue);
@@ -0,0 +1,56 @@
+////////////////////////////////////////////////////////////////////////////////
+//
+// The University of Illinois/NCSA
+// Open Source License (NCSA)
+//
+// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
+//
+// Developed by:
+//
+//                 AMD Research and AMD HSA Software Development
+//
+//                 Advanced Micro Devices, Inc.
+//
+//                 www.amd.com
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to
+// deal with the Software without restriction, including without limitation
+// the rights to use, copy, modify, merge, publish, distribute, sublicense,
+// and/or sell copies of the Software, and to permit persons to whom the
+// Software is furnished to do so, subject to the following conditions:
+//
+//  - Redistributions of source code must retain the above copyright notice,
+//    this list of conditions and the following disclaimers.
+//  - Redistributions in binary form must reproduce the above copyright
+//    notice, this list of conditions and the following disclaimers in
+//    the documentation and/or other materials provided with the distribution.
+//  - Neither the names of Advanced Micro Devices, Inc,
+//    nor the names of its contributors may be used to endorse or promote
+//    products derived from this Software without specific prior written
+//    permission.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
+// THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
+// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
+// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
+// DEALINGS WITH THE SOFTWARE.
+//
+////////////////////////////////////////////////////////////////////////////////
+
+#include "core/inc/queue.h"
+#include "core/inc/runtime.h"
+
+namespace core {
+
+void Queue::DefaultErrorHandler(hsa_status_t status, hsa_queue_t* source, void* data) {
+  if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message()) {
+    const char* msg = "UNKNOWN ERROR";
+    HSA::hsa_status_string(status, &msg);
+    fprintf(stderr, "Queue at %p inactivated due to async error:\n\t%s\n", source, msg);
+  }
+}
+
+}
@@ -64,7 +64,7 @@ class Flag {
    enable_vm_fault_message_ = (var == "0") ? false : true;

    var = os::GetEnvVar("HSA_ENABLE_QUEUE_FAULT_MESSAGE");
-    enable_queue_fault_message_ = (var == "1") ? true : false;
+    enable_queue_fault_message_ = (var == "0") ? false : true;

    var = os::GetEnvVar("HSA_ENABLE_INTERRUPT");
    enable_interrupt_ = (var == "0") ? false : true;
@@ -105,6 +105,7 @@ class Flag {
  bool enable_interrupt() const { return enable_interrupt_; }

  bool enable_thread_trace() const { return enable_thread_trace_; }
+
  bool thread_trace_buff_size() const { return thread_trace_buff_size_; }

  bool enable_sdma() const { return enable_sdma_; }