From 00a2ed60f87fe11997b2f4e16d82ba81a077fd76 Mon Sep 17 00:00:00 2001 From: Rhys Perry Date: Wed, 4 Jun 2025 16:50:25 +0100 Subject: [PATCH] radv/meta: use unsigned min in copy/fill shaders Otherwise, this would break >2 GiB copy/fill. Signed-off-by: Rhys Perry Backport: 25.1 Reviewed-by: Samuel Pitoiset Part-of: --- src/amd/vulkan/nir/radv_meta_nir.c | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/amd/vulkan/nir/radv_meta_nir.c b/src/amd/vulkan/nir/radv_meta_nir.c index 5a55007a963..0b60db3c95c 100644 --- a/src/amd/vulkan/nir/radv_meta_nir.c +++ b/src/amd/vulkan/nir/radv_meta_nir.c @@ -139,7 +139,7 @@ radv_meta_nir_build_fill_memory_shader(struct radv_device *dev, uint32_t bytes_p nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b), 0), b.shader->info.workgroup_size[0]), nir_load_local_invocation_index(&b)); - nir_def *offset = nir_imin(&b, nir_imul_imm(&b, global_id, bytes_per_invocation), max_offset); + nir_def *offset = nir_umin(&b, nir_imul_imm(&b, global_id, bytes_per_invocation), max_offset); nir_def *dst_addr = nir_iadd(&b, buffer_addr, nir_u2u64(&b, offset)); nir_build_store_global(&b, data, dst_addr, .align_mul = 4); @@ -166,7 +166,7 @@ radv_meta_nir_build_copy_memory_shader(struct radv_device *dev, uint32_t bytes_p nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b), 0), b.shader->info.workgroup_size[0]), nir_load_local_invocation_index(&b)); - nir_def *offset = nir_u2u64(&b, nir_imin(&b, nir_imul_imm(&b, global_id, bytes_per_invocation), max_offset)); + nir_def *offset = nir_u2u64(&b, nir_umin(&b, nir_imul_imm(&b, global_id, bytes_per_invocation), max_offset)); nir_def *data = nir_build_load_global(&b, num_components, bit_size, nir_iadd(&b, src_addr, offset), .align_mul = bit_size / 8);