From 9ac269fc7df5c82360d6dbab67e76712a60b1d84 Mon Sep 17 00:00:00 2001 From: Rajnesh Kanwal Date: Wed, 5 Oct 2022 16:43:35 +0500 Subject: [PATCH] pvr: Add support to generate compute kernel to update shared regs. Signed-off-by: Rajnesh Kanwal Reviewed-by: Frank Binns Part-of: --- src/imagination/vulkan/pvr_cmd_buffer.c | 39 +++++++++++++++++++++++++ src/imagination/vulkan/pvr_private.h | 15 ++++++++++ 2 files changed, 54 insertions(+) diff --git a/src/imagination/vulkan/pvr_cmd_buffer.c b/src/imagination/vulkan/pvr_cmd_buffer.c index 68712fdec3f..223e4cd0f7d 100644 --- a/src/imagination/vulkan/pvr_cmd_buffer.c +++ b/src/imagination/vulkan/pvr_cmd_buffer.c @@ -2921,6 +2921,45 @@ static void pvr_compute_update_shared(struct pvr_cmd_buffer *cmd_buffer, pvr_compute_generate_control_stream(csb, sub_cmd, &info); } +void pvr_compute_update_shared_private( + struct pvr_cmd_buffer *cmd_buffer, + struct pvr_sub_cmd_compute *const sub_cmd, + struct pvr_private_compute_pipeline *pipeline) +{ + const struct pvr_physical_device *pdevice = cmd_buffer->device->pdevice; + const uint32_t const_shared_regs = pipeline->const_shared_regs_count; + struct pvr_csb *csb = &sub_cmd->control_stream; + struct pvr_compute_kernel_info info; + + /* No shared regs, no need to use an allocation kernel. */ + if (!const_shared_regs) + return; + + info = (struct pvr_compute_kernel_info){ + .indirect_buffer_addr = PVR_DEV_ADDR_INVALID, + .usc_common_size = + DIV_ROUND_UP(const_shared_regs, + PVRX(CDMCTRL_KERNEL0_USC_COMMON_SIZE_UNIT_SIZE)), + .pds_data_size = + DIV_ROUND_UP(pipeline->pds_shared_update_data_size_dw << 2U, + PVRX(CDMCTRL_KERNEL0_PDS_DATA_SIZE_UNIT_SIZE)), + .usc_target = PVRX(CDMCTRL_USC_TARGET_ALL), + .pds_data_offset = pipeline->pds_shared_update_data_offset, + .pds_code_offset = pipeline->pds_shared_update_code_offset, + .sd_type = PVRX(CDMCTRL_SD_TYPE_NONE), + .usc_common_shared = true, + .local_size = { 1, 1, 1 }, + .global_size = { 1, 1, 1 }, + }; + + /* We don't need to pad the workgroup size. */ + + info.max_instances = + pvr_compute_flat_slot_size(pdevice, const_shared_regs, false, 1U); + + pvr_compute_generate_control_stream(csb, sub_cmd, &info); +} + static uint32_t pvr_compute_flat_pad_workgroup_size(const struct pvr_physical_device *pdevice, uint32_t workgroup_size, diff --git a/src/imagination/vulkan/pvr_private.h b/src/imagination/vulkan/pvr_private.h index 7b217313866..518ff860899 100644 --- a/src/imagination/vulkan/pvr_private.h +++ b/src/imagination/vulkan/pvr_private.h @@ -1232,6 +1232,16 @@ struct pvr_query_pool { struct pvr_bo *availability_buffer; }; +struct pvr_private_compute_pipeline { + /* Used by pvr_compute_update_shared_private(). */ + uint32_t pds_shared_update_code_offset; + uint32_t pds_shared_update_data_offset; + uint32_t pds_shared_update_data_size_dw; + uint32_t const_shared_regs_count; + + pvr_dev_addr_t const_buffer_addr; +}; + struct pvr_render_target { struct pvr_rt_dataset *rt_dataset; @@ -1534,6 +1544,11 @@ VkResult pvr_device_tile_buffer_ensure_cap(struct pvr_device *device, uint32_t capacity, uint32_t size_in_bytes); +void pvr_compute_update_shared_private( + struct pvr_cmd_buffer *cmd_buffer, + struct pvr_sub_cmd_compute *const sub_cmd, + struct pvr_private_compute_pipeline *pipeline); + #define PVR_FROM_HANDLE(__pvr_type, __name, __handle) \ VK_FROM_HANDLE(__pvr_type, __name, __handle)