From 1199f91a2f6e924e859b95b2457a2ac4c6f0e1df Mon Sep 17 00:00:00 2001 From: Samuel Pitoiset Date: Wed, 14 Jan 2026 13:05:03 +0100 Subject: [PATCH] radv/meta: use 2D array for color resolves with compute Cc: mesa-stable Signed-off-by: Samuel Pitoiset Part-of: --- src/amd/ci/radv-hawaii-fails.txt | 22 ---------------------- src/amd/ci/radv-polaris10-fails.txt | 23 ----------------------- src/amd/ci/radv-stoney-fails.txt | 5 ----- src/amd/vulkan/nir/radv_meta_nir.c | 23 +++++++++++++---------- 4 files changed, 13 insertions(+), 60 deletions(-) diff --git a/src/amd/ci/radv-hawaii-fails.txt b/src/amd/ci/radv-hawaii-fails.txt index db4d685d847..e69de29bb2d 100644 --- a/src/amd/ci/radv-hawaii-fails.txt +++ b/src/amd/ci/radv-hawaii-fails.txt @@ -1,22 +0,0 @@ -# RADV bug -dEQP-VK.pipeline.fast_linked_library.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_2,Fail -dEQP-VK.pipeline.fast_linked_library.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_4,Fail -dEQP-VK.pipeline.fast_linked_library.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_8,Fail -dEQP-VK.pipeline.monolithic.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_2,Fail -dEQP-VK.pipeline.monolithic.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_4,Fail -dEQP-VK.pipeline.monolithic.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_8,Fail -dEQP-VK.pipeline.pipeline_library.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_2,Fail -dEQP-VK.pipeline.pipeline_library.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_4,Fail -dEQP-VK.pipeline.pipeline_library.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_8,Fail -dEQP-VK.pipeline.shader_object_linked_binary.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_2,Fail -dEQP-VK.pipeline.shader_object_linked_binary.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_4,Fail -dEQP-VK.pipeline.shader_object_linked_binary.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_8,Fail -dEQP-VK.pipeline.shader_object_linked_spirv.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_2,Fail -dEQP-VK.pipeline.shader_object_linked_spirv.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_4,Fail -dEQP-VK.pipeline.shader_object_linked_spirv.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_8,Fail -dEQP-VK.pipeline.shader_object_unlinked_binary.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_2,Fail -dEQP-VK.pipeline.shader_object_unlinked_binary.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_4,Fail -dEQP-VK.pipeline.shader_object_unlinked_binary.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_8,Fail -dEQP-VK.pipeline.shader_object_unlinked_spirv.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_2,Fail -dEQP-VK.pipeline.shader_object_unlinked_spirv.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_4,Fail -dEQP-VK.pipeline.shader_object_unlinked_spirv.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_8,Fail diff --git a/src/amd/ci/radv-polaris10-fails.txt b/src/amd/ci/radv-polaris10-fails.txt index c6c1fe4691c..3575c65392c 100644 --- a/src/amd/ci/radv-polaris10-fails.txt +++ b/src/amd/ci/radv-polaris10-fails.txt @@ -4,29 +4,6 @@ dEQP-VK.video.synchronization.decode_h265.basic.event.device_set_reset,Fail dEQP-VK.video.synchronization2.decode_h264.basic.event.device_set_reset,Fail dEQP-VK.video.synchronization2.decode_h265.basic.event.device_set_reset,Fail -# RADV bug -dEQP-VK.pipeline.fast_linked_library.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_2,Fail -dEQP-VK.pipeline.fast_linked_library.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_4,Fail -dEQP-VK.pipeline.fast_linked_library.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_8,Fail -dEQP-VK.pipeline.monolithic.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_2,Fail -dEQP-VK.pipeline.monolithic.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_4,Fail -dEQP-VK.pipeline.monolithic.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_8,Fail -dEQP-VK.pipeline.pipeline_library.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_2,Fail -dEQP-VK.pipeline.pipeline_library.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_4,Fail -dEQP-VK.pipeline.pipeline_library.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_8,Fail -dEQP-VK.pipeline.shader_object_linked_binary.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_2,Fail -dEQP-VK.pipeline.shader_object_linked_binary.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_4,Fail -dEQP-VK.pipeline.shader_object_linked_binary.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_8,Fail -dEQP-VK.pipeline.shader_object_linked_spirv.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_2,Fail -dEQP-VK.pipeline.shader_object_linked_spirv.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_4,Fail -dEQP-VK.pipeline.shader_object_linked_spirv.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_8,Fail -dEQP-VK.pipeline.shader_object_unlinked_binary.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_2,Fail -dEQP-VK.pipeline.shader_object_unlinked_binary.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_4,Fail -dEQP-VK.pipeline.shader_object_unlinked_binary.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_8,Fail -dEQP-VK.pipeline.shader_object_unlinked_spirv.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_2,Fail -dEQP-VK.pipeline.shader_object_unlinked_spirv.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_4,Fail -dEQP-VK.pipeline.shader_object_unlinked_spirv.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_8,Fail - # https://gitlab.khronos.org/Tracker/vk-gl-cts/-/issues/6256 dEQP-VK.sparse_resources.image_sparse_residency.2d.r10x6_unorm_pack16.1024_128_1,Crash dEQP-VK.sparse_resources.image_sparse_residency.2d.r10x6_unorm_pack16.11_137_1,Crash diff --git a/src/amd/ci/radv-stoney-fails.txt b/src/amd/ci/radv-stoney-fails.txt index 3ba268c6ade..0ecc844274d 100644 --- a/src/amd/ci/radv-stoney-fails.txt +++ b/src/amd/ci/radv-stoney-fails.txt @@ -1,7 +1,2 @@ # See https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9049 dEQP-VK.pipeline.timestamp.calibrated.calibration_test,Fail - -# RADV bug -dEQP-VK.pipeline.pipeline_library.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_2,Fail - -dEQP-VK.pipeline.shader_object_linked_spirv.multisample.3d.64x64x8_1.r8g8b8a8_unorm.samples_4,Fail diff --git a/src/amd/vulkan/nir/radv_meta_nir.c b/src/amd/vulkan/nir/radv_meta_nir.c index 28b6bb45378..20a68b4e5db 100644 --- a/src/amd/vulkan/nir/radv_meta_nir.c +++ b/src/amd/vulkan/nir/radv_meta_nir.c @@ -1307,8 +1307,8 @@ radv_meta_nir_build_resolve_compute_shader(struct radv_device *dev, enum radv_me int samples) { enum glsl_base_type img_base_type = type == RADV_META_RESOLVE_COMPUTE_INTEGER ? GLSL_TYPE_UINT : GLSL_TYPE_FLOAT; - const struct glsl_type *sampler_type = glsl_sampler_type(GLSL_SAMPLER_DIM_MS, false, false, img_base_type); - const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_2D, false, img_base_type); + const struct glsl_type *sampler_type = glsl_sampler_type(GLSL_SAMPLER_DIM_MS, false, true, img_base_type); + const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_2D, true, img_base_type); nir_builder b = radv_meta_nir_init_shader(dev, MESA_SHADER_COMPUTE, "meta_resolve_cs-%d-%s", samples, radv_meta_resolve_compute_type_name(type)); b.shader->info.workgroup_size[0] = 8; @@ -1322,18 +1322,21 @@ radv_meta_nir_build_resolve_compute_shader(struct radv_device *dev, enum radv_me output_img->data.descriptor_set = 0; output_img->data.binding = 1; - nir_def *global_id = radv_meta_nir_get_global_ids(&b, 2); + nir_def *global_id = radv_meta_nir_get_global_ids(&b, 3); nir_def *src_offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 8); nir_def *dst_offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 8), .range = 16); - nir_def *src_coord = nir_iadd(&b, global_id, src_offset); - nir_def *dst_coord = nir_iadd(&b, global_id, dst_offset); + nir_def *src_coord = nir_iadd(&b, nir_trim_vector(&b, global_id, 2), src_offset); + nir_def *dst_coord = nir_iadd(&b, nir_trim_vector(&b, global_id, 2), dst_offset); + + nir_def *src_img_coord = + nir_vec3(&b, nir_channel(&b, src_coord, 0), nir_channel(&b, src_coord, 1), nir_channel(&b, global_id, 2)); nir_variable *color = nir_local_variable_create(b.impl, glsl_vec4_type(), "color"); radv_meta_nir_build_resolve_shader_core(dev, &b, type == RADV_META_RESOLVE_COMPUTE_INTEGER, samples, input_img, - color, src_coord); + color, src_img_coord); nir_def *outval = nir_load_var(&b, color); if (type == RADV_META_RESOLVE_COMPUTE_NORM_SRGB) @@ -1342,11 +1345,11 @@ radv_meta_nir_build_resolve_compute_shader(struct radv_device *dev, enum radv_me if (type == RADV_META_RESOLVE_COMPUTE_NORM || type == RADV_META_RESOLVE_COMPUTE_NORM_SRGB) outval = nir_f2f32(&b, nir_f2f16_rtz(&b, outval)); - nir_def *img_coord = nir_vec4(&b, nir_channel(&b, dst_coord, 0), nir_channel(&b, dst_coord, 1), nir_undef(&b, 1, 32), - nir_undef(&b, 1, 32)); + nir_def *dst_img_coord = nir_vec4(&b, nir_channel(&b, dst_coord, 0), nir_channel(&b, dst_coord, 1), + nir_channel(&b, global_id, 2), nir_undef(&b, 1, 32)); - nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->def, img_coord, nir_undef(&b, 1, 32), outval, - nir_imm_int(&b, 0), .image_dim = GLSL_SAMPLER_DIM_2D); + nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->def, dst_img_coord, nir_undef(&b, 1, 32), outval, + nir_imm_int(&b, 0), .image_dim = GLSL_SAMPLER_DIM_2D, .image_array = true); return b.shader; }