From 5e4523348408c989d50a4b31582bdb28e03bb599 Mon Sep 17 00:00:00 2001 From: Matthew Poremba Date: Sat, 11 Nov 2023 13:01:44 -0600 Subject: [PATCH] gpu-compute: Add gfx version to HSA task entry The version is necessary for determining the correct ABI init process. Add it to the task queue so it is accessible when doing ABI init. Change-Id: If77434b0f93614057b5c40fcf612d59b54e05dbb --- src/gpu-compute/hsa_queue_entry.hh | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) diff --git a/src/gpu-compute/hsa_queue_entry.hh b/src/gpu-compute/hsa_queue_entry.hh index 84ae139127..7e6744704a 100644 --- a/src/gpu-compute/hsa_queue_entry.hh +++ b/src/gpu-compute/hsa_queue_entry.hh @@ -63,7 +63,7 @@ class HSAQueueEntry HSAQueueEntry(std::string kernel_name, uint32_t queue_id, int dispatch_id, void *disp_pkt, AMDKernelCode *akc, Addr host_pkt_addr, Addr code_addr, GfxVersion gfx_version) - : kernName(kernel_name), + : _gfxVersion(gfx_version), kernName(kernel_name), _wgSize{{(int)((_hsa_dispatch_packet_t*)disp_pkt)->workgroup_size_x, (int)((_hsa_dispatch_packet_t*)disp_pkt)->workgroup_size_y, (int)((_hsa_dispatch_packet_t*)disp_pkt)->workgroup_size_z}}, @@ -129,6 +129,12 @@ class HSAQueueEntry parseKernelCode(akc); } + const GfxVersion& + gfxVersion() const + { + return _gfxVersion; + } + const std::string& kernelName() const { @@ -438,6 +444,8 @@ class HSAQueueEntry initialVgprState.set(WorkitemIdZ, akc->enable_vgpr_workitem_id > 1); } + // store gfx version for version specific task handling + GfxVersion _gfxVersion; // name of the kernel associated with the AQL entry std::string kernName; // workgroup Size (3 dimensions)