From c34655f1ffeb4fb06d9b0667cdd99b4d9858ae8f Mon Sep 17 00:00:00 2001 From: Evgeny Date: Fri, 8 Nov 2019 18:08:09 -0600 Subject: [PATCH] setting of granulated lds, sgpr, vgpr reserved by HW for the kernel Change-Id: Ib3e60c0b97787f56f769136480a46a8560e00cc8 [ROCm/rocprofiler commit: dc05345a78e0a9c649d9f1a27c056aa065453b5e] --- projects/rocprofiler/test/tool/tool.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/rocprofiler/test/tool/tool.cpp b/projects/rocprofiler/test/tool/tool.cpp index 877fd5b984..1b0ad02abc 100644 --- a/projects/rocprofiler/test/tool/tool.cpp +++ b/projects/rocprofiler/test/tool/tool.cpp @@ -658,10 +658,10 @@ hsa_status_t dispatch_callback(const rocprofiler_callback_data_t* callback_data, uint64_t workgroup_size = packet->workgroup_size_x * packet->workgroup_size_y * packet->workgroup_size_z; if (workgroup_size > UINT32_MAX) abort(); kernel_properties_ptr->workgroup_size = (uint32_t)workgroup_size; - kernel_properties_ptr->lds_size = packet->group_segment_size; + kernel_properties_ptr->lds_size = AMD_HSA_BITS_GET(kernel_code->compute_pgm_rsrc2, AMD_COMPUTE_PGM_RSRC_TWO_GRANULATED_LDS_SIZE); // packet->group_segment_size; kernel_properties_ptr->scratch_size = packet->private_segment_size; - kernel_properties_ptr->vgpr_count = kernel_code->reserved_vgpr_count; - kernel_properties_ptr->sgpr_count = kernel_code->reserved_sgpr_count; + kernel_properties_ptr->vgpr_count = AMD_HSA_BITS_GET(kernel_code->compute_pgm_rsrc1, AMD_COMPUTE_PGM_RSRC_ONE_GRANULATED_WORKITEM_VGPR_COUNT); + kernel_properties_ptr->sgpr_count = AMD_HSA_BITS_GET(kernel_code->compute_pgm_rsrc1, AMD_COMPUTE_PGM_RSRC_ONE_GRANULATED_WAVEFRONT_SGPR_COUNT); kernel_properties_ptr->fbarrier_count = kernel_code->workgroup_fbarrier_count; kernel_properties_ptr->signal = callback_data->completion_signal;