From 297cf6f1aa8019955d8abd52f154b5f1f7163888 Mon Sep 17 00:00:00 2001 From: Samuel Pitoiset Date: Wed, 16 Jul 2025 09:56:30 +0200 Subject: [PATCH] radv/meta: add a pass to clear HiZ surfaces Signed-off-by: Samuel Pitoiset Part-of: --- src/amd/vulkan/meson.build | 1 + src/amd/vulkan/meta/radv_meta.h | 4 + src/amd/vulkan/meta/radv_meta_clear_hiz.c | 184 ++++++++++++++++++++++ src/amd/vulkan/nir/radv_meta_nir.c | 34 ++++ src/amd/vulkan/nir/radv_meta_nir.h | 2 + 5 files changed, 225 insertions(+) create mode 100644 src/amd/vulkan/meta/radv_meta_clear_hiz.c diff --git a/src/amd/vulkan/meson.build b/src/amd/vulkan/meson.build index 158bc4d7c03..e1974261803 100644 --- a/src/amd/vulkan/meson.build +++ b/src/amd/vulkan/meson.build @@ -51,6 +51,7 @@ libradv_files = files( 'meta/radv_meta_buffer.c', 'meta/radv_meta_bufimage.c', 'meta/radv_meta_clear.c', + 'meta/radv_meta_clear_hiz.c', 'meta/radv_meta_copy.c', 'meta/radv_meta_copy_vrs_htile.c', 'meta/radv_meta_dcc_retile.c', diff --git a/src/amd/vulkan/meta/radv_meta.h b/src/amd/vulkan/meta/radv_meta.h index be8005f0587..99f906f442b 100644 --- a/src/amd/vulkan/meta/radv_meta.h +++ b/src/amd/vulkan/meta/radv_meta.h @@ -93,6 +93,7 @@ enum radv_meta_object_key_type { RADV_META_OBJECT_KEY_CLEAR_DS, RADV_META_OBJECT_KEY_CLEAR_HTILE, RADV_META_OBJECT_KEY_CLEAR_DCC_COMP_TO_SINGLE, + RADV_META_OBJECT_KEY_CLEAR_HIZ, RADV_META_OBJECT_KEY_FAST_CLEAR_ELIMINATE, RADV_META_OBJECT_KEY_DCC_DECOMPRESS, RADV_META_OBJECT_KEY_DCC_RETILE, @@ -240,6 +241,9 @@ uint32_t radv_clear_dcc(struct radv_cmd_buffer *cmd_buffer, struct radv_image *i uint32_t radv_clear_htile(struct radv_cmd_buffer *cmd_buffer, const struct radv_image *image, const VkImageSubresourceRange *range, uint32_t value, bool is_clear); +void radv_clear_hiz(struct radv_cmd_buffer *cmd_buffer, struct radv_image *image, const VkImageSubresourceRange *range, + uint32_t value); + void radv_update_memory_cp(struct radv_cmd_buffer *cmd_buffer, uint64_t va, const void *data, uint64_t size); void radv_update_memory(struct radv_cmd_buffer *cmd_buffer, uint64_t va, uint64_t size, const void *data, diff --git a/src/amd/vulkan/meta/radv_meta_clear_hiz.c b/src/amd/vulkan/meta/radv_meta_clear_hiz.c new file mode 100644 index 00000000000..5a6400b6e4a --- /dev/null +++ b/src/amd/vulkan/meta/radv_meta_clear_hiz.c @@ -0,0 +1,184 @@ +/* + * Copyright © 2025 Valve Corporation + * + * SPDX-License-Identifier: MIT + */ + +#include "nir/radv_meta_nir.h" +#include "radv_entrypoints.h" +#include "radv_meta.h" + +static VkResult +get_clear_hiz_pipeline_layout(struct radv_device *device, VkPipelineLayout *layout_out) +{ + enum radv_meta_object_key_type key = RADV_META_OBJECT_KEY_CLEAR_HIZ; + + const VkDescriptorSetLayoutBinding binding = { + .binding = 0, + .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE, + .descriptorCount = 1, + .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT, + }; + + const VkDescriptorSetLayoutCreateInfo desc_info = { + .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO, + .flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT, + .bindingCount = 1, + .pBindings = &binding, + }; + + const VkPushConstantRange pc_range = { + .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT, + .size = 4, + }; + + return vk_meta_get_pipeline_layout(&device->vk, &device->meta_state.device, &desc_info, &pc_range, &key, sizeof(key), + layout_out); +} + +struct radv_clear_hiz_key { + enum radv_meta_object_key_type type; + uint8_t samples; +}; + +static VkResult +get_clear_hiz_pipeline(struct radv_device *device, const struct radv_image *image, VkPipeline *pipeline_out, + VkPipelineLayout *layout_out) +{ + const uint32_t samples = image->vk.samples; + struct radv_clear_hiz_key key; + VkResult result; + + result = get_clear_hiz_pipeline_layout(device, layout_out); + if (result != VK_SUCCESS) + return result; + + memset(&key, 0, sizeof(key)); + key.type = RADV_META_OBJECT_KEY_CLEAR_HIZ; + key.samples = samples; + + VkPipeline pipeline_from_cache = vk_meta_lookup_pipeline(&device->meta_state.device, &key, sizeof(key)); + if (pipeline_from_cache != VK_NULL_HANDLE) { + *pipeline_out = pipeline_from_cache; + return VK_SUCCESS; + } + + nir_shader *cs = radv_meta_nir_build_clear_hiz_compute_shader(device, samples); + + const VkPipelineShaderStageCreateInfo stage_info = { + .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO, + .stage = VK_SHADER_STAGE_COMPUTE_BIT, + .module = vk_shader_module_handle_from_nir(cs), + .pName = "main", + .pSpecializationInfo = NULL, + }; + + const VkComputePipelineCreateInfo pipeline_info = { + .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO, + .stage = stage_info, + .flags = 0, + .layout = *layout_out, + }; + + result = vk_meta_create_compute_pipeline(&device->vk, &device->meta_state.device, &pipeline_info, &key, sizeof(key), + pipeline_out); + + ralloc_free(cs); + return result; +} + +void +radv_clear_hiz(struct radv_cmd_buffer *cmd_buffer, struct radv_image *image, const VkImageSubresourceRange *range, + uint32_t value) +{ + struct radv_device *device = radv_cmd_buffer_device(cmd_buffer); + const struct radeon_surf *surf = &image->planes[0].surface; + struct radv_meta_saved_state saved_state; + struct radv_image_view iview; + VkPipelineLayout layout; + VkPipeline pipeline; + VkResult result; + + /* Clearing HiZ should only be needed to implement a workaround on GFX12. */ + assert(image->hiz_valid_offset); + + result = get_clear_hiz_pipeline(device, image, &pipeline, &layout); + if (result != VK_SUCCESS) { + vk_command_buffer_set_error(&cmd_buffer->vk, result); + return; + } + + cmd_buffer->state.flush_bits |= + radv_src_access_flush(cmd_buffer, VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT, + VK_ACCESS_2_DEPTH_STENCIL_ATTACHMENT_WRITE_BIT, 0, image, range); + + radv_meta_save(&saved_state, cmd_buffer, + RADV_META_SAVE_COMPUTE_PIPELINE | RADV_META_SAVE_DESCRIPTORS | RADV_META_SAVE_CONSTANTS); + + radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE, pipeline); + + const uint32_t base_width = surf->u.gfx9.zs.hiz.width_in_tiles; + const uint32_t base_height = surf->u.gfx9.zs.hiz.height_in_tiles; + + for (uint32_t l = 0; l < vk_image_subresource_level_count(&image->vk, range); l++) { + uint32_t width, height; + + width = u_minify(base_width, range->baseMipLevel + l); + height = u_minify(base_height, range->baseMipLevel + l); + + for (uint32_t s = 0; s < vk_image_subresource_layer_count(&image->vk, range); s++) { + radv_hiz_image_view_init(&iview, device, + &(VkImageViewCreateInfo){ + .sType = VK_STRUCTURE_TYPE_IMAGE_VIEW_CREATE_INFO, + .image = radv_image_to_handle(image), + .viewType = radv_meta_get_view_type(image), + .format = image->vk.format, + .subresourceRange = + { + .aspectMask = VK_IMAGE_ASPECT_DEPTH_BIT, + .baseMipLevel = range->baseMipLevel + l, + .levelCount = 1, + .baseArrayLayer = range->baseArrayLayer + s, + .layerCount = 1, + }, + }); + + radv_meta_bind_descriptors(cmd_buffer, VK_PIPELINE_BIND_POINT_COMPUTE, layout, 1, + (VkDescriptorGetInfoEXT[]){ + { + .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_GET_INFO_EXT, + .type = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE, + .data.pStorageImage = + (VkDescriptorImageInfo[]){ + { + .sampler = VK_NULL_HANDLE, + .imageView = radv_image_view_to_handle(&iview), + .imageLayout = VK_IMAGE_LAYOUT_GENERAL, + }, + }, + }, + }); + + const VkPushConstantsInfo pc_info = { + .sType = VK_STRUCTURE_TYPE_PUSH_CONSTANTS_INFO_KHR, + .layout = layout, + .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT, + .offset = 0, + .size = sizeof(value), + .pValues = &value, + }; + + radv_CmdPushConstants2(radv_cmd_buffer_to_handle(cmd_buffer), &pc_info); + + radv_unaligned_dispatch(cmd_buffer, width, height, 1); + + radv_image_view_finish(&iview); + } + } + + radv_meta_restore(&saved_state, cmd_buffer); + + cmd_buffer->state.flush_bits |= + RADV_CMD_FLAG_CS_PARTIAL_FLUSH | radv_src_access_flush(cmd_buffer, VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT, + VK_ACCESS_2_SHADER_WRITE_BIT, 0, image, range); +} diff --git a/src/amd/vulkan/nir/radv_meta_nir.c b/src/amd/vulkan/nir/radv_meta_nir.c index e76c3a3aade..1e367817d89 100644 --- a/src/amd/vulkan/nir/radv_meta_nir.c +++ b/src/amd/vulkan/nir/radv_meta_nir.c @@ -1527,3 +1527,37 @@ radv_meta_nir_build_resolve_fs(struct radv_device *dev) return b.shader; } + +nir_shader * +radv_meta_nir_build_clear_hiz_compute_shader(struct radv_device *dev, int samples) +{ + const enum glsl_sampler_dim dim = samples > 1 ? GLSL_SAMPLER_DIM_MS : GLSL_SAMPLER_DIM_2D; + const struct glsl_type *img_type = glsl_image_type(dim, false, GLSL_TYPE_FLOAT); + nir_builder b = radv_meta_nir_init_shader(dev, MESA_SHADER_COMPUTE, "meta_clear_hiz_cs-%d", samples); + b.shader->info.workgroup_size[0] = 8; + b.shader->info.workgroup_size[1] = 8; + + nir_variable *output_img = nir_variable_create(b.shader, nir_var_image, img_type, "out_img"); + output_img->data.descriptor_set = 0; + output_img->data.binding = 0; + + nir_def *global_id = radv_meta_nir_get_global_ids(&b, 2); + + nir_def *clear_val = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .range = 4); + + nir_def *comps[4]; + comps[0] = nir_channel(&b, global_id, 0); + comps[1] = nir_channel(&b, global_id, 1); + comps[2] = nir_imm_int(&b, 0); + comps[3] = nir_undef(&b, 1, 32); + global_id = nir_vec(&b, comps, 4); + + nir_def *data = nir_vec4(&b, clear_val, nir_imm_int(&b, 0), nir_imm_int(&b, 0), nir_imm_int(&b, 0)); + + for (uint32_t i = 0; i < samples; i++) { + nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->def, global_id, nir_imm_int(&b, i), data, + nir_imm_int(&b, 0), .image_dim = dim); + } + + return b.shader; +} diff --git a/src/amd/vulkan/nir/radv_meta_nir.h b/src/amd/vulkan/nir/radv_meta_nir.h index 43263326f34..b193eee18e1 100644 --- a/src/amd/vulkan/nir/radv_meta_nir.h +++ b/src/amd/vulkan/nir/radv_meta_nir.h @@ -109,6 +109,8 @@ nir_shader *radv_meta_nir_build_depth_stencil_resolve_fragment_shader(struct rad nir_shader *radv_meta_nir_build_resolve_fs(struct radv_device *dev); +nir_shader *radv_meta_nir_build_clear_hiz_compute_shader(struct radv_device *dev, int samples); + #ifdef __cplusplus } #endif