mirror of
https://gitlab.freedesktop.org/mesa/mesa.git
synced 2026-05-06 07:18:17 +02:00
radv/meta: add a pass to clear HiZ surfaces
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36739>
This commit is contained in:
Samuel Pitoiset
Marge Bot
parent
8886a3385b
commit
297cf6f1aa
5 changed files with 225 additions and 0 deletions
|
|
@ -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',
|
||||
|
|
|
|||
|
|
@ -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,
|
||||
|
|
|
|||
184
src/amd/vulkan/meta/radv_meta_clear_hiz.c
Normal file
184
src/amd/vulkan/meta/radv_meta_clear_hiz.c
Normal file
|
|
@ -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);
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
Loading…
Add table
Reference in a new issue