From d657afa812e5bc1c071bfab4cff0e3892181294e Mon Sep 17 00:00:00 2001 From: Evgeny Date: Tue, 18 Feb 2020 14:00:19 -0600 Subject: [PATCH] adding AgentInfo::lds_block_size Change-Id: I186893add96dc92570e710ae78b475897ebfe531 [ROCm/rocprofiler commit: 40730e34e41cba43d139db97c485370890875b3d] --- projects/rocprofiler/src/util/hsa_rsrc_factory.h | 3 ++- projects/rocprofiler/test/tool/tool.cpp | 2 +- projects/rocprofiler/test/util/hsa_rsrc_factory.h | 3 ++- 3 files changed, 5 insertions(+), 3 deletions(-) diff --git a/projects/rocprofiler/src/util/hsa_rsrc_factory.h b/projects/rocprofiler/src/util/hsa_rsrc_factory.h index a6a3f36c3e..893d7f66b5 100644 --- a/projects/rocprofiler/src/util/hsa_rsrc_factory.h +++ b/projects/rocprofiler/src/util/hsa_rsrc_factory.h @@ -164,10 +164,11 @@ struct AgentInfo { // Number of Shader Arrays Per Shader Engines in Gpu uint32_t shader_arrays_per_se; - // SGPR/VGPR block sizes + // SGPR/VGPR/LDS block sizes uint32_t sgpr_block_dflt; uint32_t sgpr_block_size; uint32_t vgpr_block_size; + static const uint32_t lds_block_size = 128 * 4; }; // HSA timer class diff --git a/projects/rocprofiler/test/tool/tool.cpp b/projects/rocprofiler/test/tool/tool.cpp index 463590858d..a429c493e6 100644 --- a/projects/rocprofiler/test/tool/tool.cpp +++ b/projects/rocprofiler/test/tool/tool.cpp @@ -483,7 +483,7 @@ bool dump_context_entry(context_entry_t* entry) { entry->data.thread_id, entry->kernel_properties.grid_size, entry->kernel_properties.workgroup_size, - (entry->kernel_properties.lds_size * (128 * 4)), + entry->kernel_properties.lds_size * AgentInfo::lds_block_size, entry->kernel_properties.scratch_size, (entry->kernel_properties.vgpr_count + 1) * agent_info->vgpr_block_size, (entry->kernel_properties.sgpr_count + agent_info->sgpr_block_dflt) * agent_info->sgpr_block_size, diff --git a/projects/rocprofiler/test/util/hsa_rsrc_factory.h b/projects/rocprofiler/test/util/hsa_rsrc_factory.h index b075219fe8..3108463b20 100644 --- a/projects/rocprofiler/test/util/hsa_rsrc_factory.h +++ b/projects/rocprofiler/test/util/hsa_rsrc_factory.h @@ -162,10 +162,11 @@ struct AgentInfo { // Number of Shader Arrays Per Shader Engines in Gpu uint32_t shader_arrays_per_se; - // SGPR/VGPR block sizes + // SGPR/VGPR/LDS block sizes uint32_t sgpr_block_dflt; uint32_t sgpr_block_size; uint32_t vgpr_block_size; + static const uint32_t lds_block_size = 128 * 4; }; // HSA timer class