fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-08 04:48:08 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

radv/meta: use 2D array for color resolves with compute

Browse source 
Cc: mesa-stable
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/39285>
This commit is contained in:
Samuel Pitoiset 2026-01-14 13:05:03 +01:00 • committed by Marge Bot
parent 939eab79bb
commit 1199f91a2f
4 changed files with 13 additions and 60 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

22
src/amd/ci/radv-hawaii-fails.txt
View file

@ -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

23
src/amd/ci/radv-polaris10-fails.txt
View file

@ -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

5
src/amd/ci/radv-stoney-fails.txt
View file

@ -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

23
src/amd/vulkan/nir/radv_meta_nir.c
View file

@ -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;
}