diff --git a/projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt b/projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt
index 5561d0cf2b..00ac394805 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt
+++ b/projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt
@@ -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"
diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h
index b73294faf8..5ec43bd7b9 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h
@@ -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_;
 
diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp
index 3b01d2bc97..1d1633f66e 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp
@@ -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) {
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 e698099d18..9fcf50bf73 100755
--- 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
@@ -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) {
diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_memory_region.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_memory_region.cpp
index 95b4f259ab..295be378d9 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_memory_region.cpp
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_memory_region.cpp
@@ -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:
diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp
index 66927b5fe3..08852b9b40 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp
@@ -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);
diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/queue.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/queue.cpp
new file mode 100644
index 0000000000..9246d4a04d
--- /dev/null
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/queue.cpp
@@ -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);
+  }
+}
+
+}
diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/util/flag.h b/projects/rocr-runtime/runtime/hsa-runtime/core/util/flag.h
index 1e1fe0f833..d7add470c5 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/util/flag.h
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/util/flag.h
@@ -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_; }