diff --git a/src/amd/vulkan/radv_meta_clear.c b/src/amd/vulkan/radv_meta_clear.c index 21d4ab199e4..2fd99c0dbb6 100644 --- a/src/amd/vulkan/radv_meta_clear.c +++ b/src/amd/vulkan/radv_meta_clear.c @@ -318,8 +318,10 @@ finish_meta_clear_dcc_comp_to_single_state(struct radv_device *device) { struct radv_meta_state *state = &device->meta_state; - radv_DestroyPipeline(radv_device_to_handle(device), state->clear_dcc_comp_to_single_pipeline, - &state->alloc); + for (uint32_t i = 0; i < 2; i++) { + radv_DestroyPipeline(radv_device_to_handle(device), + state->clear_dcc_comp_to_single_pipeline[i], &state->alloc); + } radv_DestroyPipelineLayout(radv_device_to_handle(device), state->clear_dcc_comp_to_single_p_layout, &state->alloc); radv_DestroyDescriptorSetLayout(radv_device_to_handle(device), state->clear_dcc_comp_to_single_ds_layout, @@ -1147,13 +1149,18 @@ fail: return result; } +/* Clear DCC using comp-to-single by storing the clear value at the beginning of every 256B block. + * For MSAA images, clearing the first sample should be enough as long as CMASK is also cleared. + */ static nir_shader * -build_clear_dcc_comp_to_single_shader() +build_clear_dcc_comp_to_single_shader(bool is_msaa) { - const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_2D, true, GLSL_TYPE_FLOAT); + enum glsl_sampler_dim dim = is_msaa ? GLSL_SAMPLER_DIM_MS : GLSL_SAMPLER_DIM_2D; + const struct glsl_type *img_type = glsl_image_type(dim, true, GLSL_TYPE_FLOAT); nir_builder b = - nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "meta_clear_dcc_comp_to_single"); + nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "meta_clear_dcc_comp_to_single-%s", + is_msaa ? "multisampled" : "singlesampled"); b.shader->info.workgroup_size[0] = 8; b.shader->info.workgroup_size[1] = 8; b.shader->info.workgroup_size[2] = 1; @@ -1190,19 +1197,20 @@ build_clear_dcc_comp_to_single_shader() nir_channel(&b, clear_values, 1)); /* Store the clear color values. */ + nir_ssa_def *sample_id = is_msaa ? nir_imm_int(&b, 0) : nir_ssa_undef(&b, 1, 32); nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->dest.ssa, coord, - nir_imm_int(&b, 0), data, nir_imm_int(&b, 0), - .image_dim = GLSL_SAMPLER_DIM_2D, .image_array = true); + sample_id, data, nir_imm_int(&b, 0), + .image_dim = dim, .image_array = true); return b.shader; } static VkResult -create_dcc_comp_to_single_pipeline(struct radv_device *device, VkPipeline *pipeline) +create_dcc_comp_to_single_pipeline(struct radv_device *device, bool is_msaa, VkPipeline *pipeline) { struct radv_meta_state *state = &device->meta_state; VkResult result; - nir_shader *cs = build_clear_dcc_comp_to_single_shader(); + nir_shader *cs = build_clear_dcc_comp_to_single_shader(is_msaa); VkPipelineShaderStageCreateInfo shader_stage = { .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO, @@ -1268,9 +1276,12 @@ init_meta_clear_dcc_comp_to_single_state(struct radv_device *device) if (result != VK_SUCCESS) goto fail; - result = create_dcc_comp_to_single_pipeline(device, &state->clear_dcc_comp_to_single_pipeline); - if (result != VK_SUCCESS) - goto fail; + for (uint32_t i = 0; i < 2; i++) { + result = create_dcc_comp_to_single_pipeline(device, !!i, + &state->clear_dcc_comp_to_single_pipeline[i]); + if (result != VK_SUCCESS) + goto fail; + } fail: return result; @@ -1513,6 +1524,7 @@ radv_clear_dcc_comp_to_single(struct radv_cmd_buffer *cmd_buffer, unsigned bytes_per_pixel = vk_format_get_blocksize(image->vk_format); unsigned layer_count = radv_get_layerCount(image, range); struct radv_meta_saved_state saved_state; + bool is_msaa = image->info.samples > 1; struct radv_image_view iview; VkFormat format; @@ -1540,7 +1552,7 @@ radv_clear_dcc_comp_to_single(struct radv_cmd_buffer *cmd_buffer, &saved_state, cmd_buffer, RADV_META_SAVE_DESCRIPTORS | RADV_META_SAVE_COMPUTE_PIPELINE | RADV_META_SAVE_CONSTANTS); - VkPipeline pipeline = device->meta_state.clear_dcc_comp_to_single_pipeline; + VkPipeline pipeline = device->meta_state.clear_dcc_comp_to_single_pipeline[is_msaa]; radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE, pipeline); diff --git a/src/amd/vulkan/radv_private.h b/src/amd/vulkan/radv_private.h index 9005f53701c..7f09157e2d3 100644 --- a/src/amd/vulkan/radv_private.h +++ b/src/amd/vulkan/radv_private.h @@ -475,7 +475,7 @@ struct radv_meta_state { VkDescriptorSetLayout copy_vrs_htile_ds_layout; /* Clear DCC with comp-to-single. */ - VkPipeline clear_dcc_comp_to_single_pipeline; + VkPipeline clear_dcc_comp_to_single_pipeline[2]; /* 0: 1x, 1: 2x/4x/8x */ VkPipelineLayout clear_dcc_comp_to_single_p_layout; VkDescriptorSetLayout clear_dcc_comp_to_single_ds_layout;