From a65f3f5b71c65d0831fc73a7ee08cc63cae618cf Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Sat, 15 Jan 2022 15:26:07 -0600 Subject: [PATCH] Improve scratch error detection in debug mode. Adds asserts for invalid dispatch dims and scratch requests that don't actually use scratch. Change-Id: I6e6eef3f17dc38adaf96550fa55bd8625868efa3 --- runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp index 486b63cd00..ac452a34e1 100644 --- a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp @@ -824,8 +824,12 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { assert(pkt.IsValid() && "Invalid packet in dynamic scratch handler."); assert(pkt.type() == HSA_PACKET_TYPE_KERNEL_DISPATCH && "Invalid packet in dynamic scratch handler."); + assert((pkt.dispatch.workgroup_size_x != 0) && (pkt.dispatch.workgroup_size_y != 0) && + (pkt.dispatch.workgroup_size_z != 0) && "Invalid dispatch dimension."); uint32_t scratch_request = pkt.dispatch.private_segment_size; + assert((scratch_request != 0) && + "Scratch memory request from packet with no scratch demand. Possible bad kernel code object."); const uint32_t MaxScratchSlots = (queue->amd_queue_.max_cu_id + 1) * queue->agent_->properties().MaxSlotsScratchCU;