From f741c04ed10a5ebaebf786ba5cd8b75581fc08a1 Mon Sep 17 00:00:00 2001 From: Samuel Pitoiset Date: Thu, 26 Aug 2021 09:02:07 +0200 Subject: [PATCH] radv: use get_global_ids() to compute coordinates in meta shaders This was duplicated everywhere. Signed-off-by: Samuel Pitoiset Reviewed-by: Bas Nieuwenhuizen Part-of: --- src/amd/vulkan/radv_meta.c | 16 +++++ src/amd/vulkan/radv_meta.h | 2 + src/amd/vulkan/radv_meta_buffer.c | 16 +---- src/amd/vulkan/radv_meta_bufimage.c | 71 +++++++---------------- src/amd/vulkan/radv_meta_clear.c | 18 +----- src/amd/vulkan/radv_meta_copy_vrs_htile.c | 13 +---- src/amd/vulkan/radv_meta_dcc_retile.c | 16 ----- src/amd/vulkan/radv_meta_fast_clear.c | 16 +++-- src/amd/vulkan/radv_meta_fmask_expand.c | 12 +--- src/amd/vulkan/radv_meta_resolve_cs.c | 32 ++++------ src/amd/vulkan/radv_query.c | 32 ++-------- 11 files changed, 70 insertions(+), 174 deletions(-) diff --git a/src/amd/vulkan/radv_meta.c b/src/amd/vulkan/radv_meta.c index 5475e373417..8e2a9180d04 100644 --- a/src/amd/vulkan/radv_meta.c +++ b/src/amd/vulkan/radv_meta.c @@ -691,3 +691,19 @@ radv_meta_load_descriptor(nir_builder *b, unsigned desc_set, unsigned binding) .binding = binding); return nir_channels(b, rsrc, 0x3); } + +nir_ssa_def * +get_global_ids(nir_builder *b, unsigned num_components) +{ + unsigned mask = BITFIELD_MASK(num_components); + + nir_ssa_def *local_ids = nir_channels(b, nir_load_local_invocation_id(b), mask); + nir_ssa_def *block_ids = nir_channels(b, nir_load_workgroup_id(b, 32), mask); + nir_ssa_def *block_size = nir_channels( + b, + nir_imm_ivec4(b, b->shader->info.workgroup_size[0], b->shader->info.workgroup_size[1], + b->shader->info.workgroup_size[2], 0), + mask); + + return nir_iadd(b, nir_imul(b, block_ids, block_size), local_ids); +} diff --git a/src/amd/vulkan/radv_meta.h b/src/amd/vulkan/radv_meta.h index 00d3311f9d0..cfc5a5faff1 100644 --- a/src/amd/vulkan/radv_meta.h +++ b/src/amd/vulkan/radv_meta.h @@ -291,6 +291,8 @@ void radv_meta_build_resolve_shader_core(nir_builder *b, bool is_integer, int sa nir_ssa_def *radv_meta_load_descriptor(nir_builder *b, unsigned desc_set, unsigned binding); +nir_ssa_def *get_global_ids(nir_builder *b, unsigned num_components); + #ifdef __cplusplus } #endif diff --git a/src/amd/vulkan/radv_meta_buffer.c b/src/amd/vulkan/radv_meta_buffer.c index 644558c50dc..b66bb57a1ac 100644 --- a/src/amd/vulkan/radv_meta_buffer.c +++ b/src/amd/vulkan/radv_meta_buffer.c @@ -12,13 +12,7 @@ build_buffer_fill_shader(struct radv_device *dev) b.shader->info.workgroup_size[1] = 1; b.shader->info.workgroup_size[2] = 1; - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); + nir_ssa_def *global_id = get_global_ids(&b, 1); nir_ssa_def *offset = nir_imul(&b, global_id, nir_imm_int(&b, 16)); offset = nir_channel(&b, offset, 0); @@ -42,13 +36,7 @@ build_buffer_copy_shader(struct radv_device *dev) b.shader->info.workgroup_size[1] = 1; b.shader->info.workgroup_size[2] = 1; - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); + nir_ssa_def *global_id = get_global_ids(&b, 1); nir_ssa_def *offset = nir_imul(&b, global_id, nir_imm_int(&b, 16)); offset = nir_channel(&b, offset, 0); diff --git a/src/amd/vulkan/radv_meta_bufimage.c b/src/amd/vulkan/radv_meta_bufimage.c index de5855e2e73..e42da97f259 100644 --- a/src/amd/vulkan/radv_meta_bufimage.c +++ b/src/amd/vulkan/radv_meta_bufimage.c @@ -51,13 +51,7 @@ build_nir_itob_compute_shader(struct radv_device *dev, bool is_3d) output_img->data.descriptor_set = 0; output_img->data.binding = 1; - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); + nir_ssa_def *global_id = get_global_ids(&b, is_3d ? 3 : 2); nir_ssa_def *offset = nir_load_push_constant(&b, is_3d ? 3 : 2, 32, nir_imm_int(&b, 0), .range = 16); @@ -239,13 +233,7 @@ build_nir_btoi_compute_shader(struct radv_device *dev, bool is_3d) output_img->data.descriptor_set = 0; output_img->data.binding = 1; - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); + nir_ssa_def *global_id = get_global_ids(&b, is_3d ? 3 : 2); nir_ssa_def *offset = nir_load_push_constant(&b, is_3d ? 3 : 2, 32, nir_imm_int(&b, 0), .range = 16); @@ -257,7 +245,7 @@ build_nir_btoi_compute_shader(struct radv_device *dev, bool is_3d) nir_ssa_def *buf_coord = nir_imul(&b, pos_y, stride); buf_coord = nir_iadd(&b, buf_coord, pos_x); - nir_ssa_def *img_coord = nir_iadd(&b, global_id, offset); + nir_ssa_def *coord = nir_iadd(&b, global_id, offset); nir_ssa_def *input_img_deref = &nir_build_deref_var(&b, input_img)->dest.ssa; nir_tex_instr *tex = nir_tex_instr_create(b.shader, 3); @@ -277,6 +265,12 @@ build_nir_btoi_compute_shader(struct radv_device *dev, bool is_3d) nir_builder_instr_insert(&b, &tex->instr); nir_ssa_def *outval = &tex->dest.ssa; + + nir_ssa_def *img_coord = nir_vec4(&b, nir_channel(&b, coord, 0), + nir_channel(&b, coord, 1), + is_3d ? nir_channel(&b, coord, 2) : nir_ssa_undef(&b, 1, 32), + nir_ssa_undef(&b, 1, 32)); + nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->dest.ssa, img_coord, nir_ssa_undef(&b, 1, 32), outval, nir_imm_int(&b, 0), .image_dim = dim); @@ -419,13 +413,7 @@ build_nir_btoi_r32g32b32_compute_shader(struct radv_device *dev) output_img->data.descriptor_set = 0; output_img->data.binding = 1; - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); + nir_ssa_def *global_id = get_global_ids(&b, 2); nir_ssa_def *offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 16); nir_ssa_def *pitch = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 8), .range = 16); @@ -579,13 +567,7 @@ build_nir_itoi_compute_shader(struct radv_device *dev, bool is_3d, int samples) output_img->data.descriptor_set = 0; output_img->data.binding = 1; - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); + nir_ssa_def *global_id = get_global_ids(&b, is_3d ? 3 : 2); nir_ssa_def *src_offset = nir_load_push_constant(&b, is_3d ? 3 : 2, 32, nir_imm_int(&b, 0), .range = 24); @@ -622,9 +604,14 @@ build_nir_itoi_compute_shader(struct radv_device *dev, bool is_3d, int samples) nir_builder_instr_insert(&b, &tex->instr); } + nir_ssa_def *img_coord = nir_vec4(&b, nir_channel(&b, dst_coord, 0), + nir_channel(&b, dst_coord, 1), + is_3d ? nir_channel(&b, dst_coord, 2) : nir_ssa_undef(&b, 1, 32), + nir_ssa_undef(&b, 1, 32)); + for (uint32_t i = 0; i < samples; i++) { nir_ssa_def *outval = &tex_instr[i]->dest.ssa; - nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->dest.ssa, dst_coord, + nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->dest.ssa, img_coord, nir_imm_int(&b, i), outval, nir_imm_int(&b, 0), .image_dim = dim); } @@ -781,13 +768,7 @@ build_nir_itoi_r32g32b32_compute_shader(struct radv_device *dev) output_img->data.descriptor_set = 0; output_img->data.binding = 1; - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); + nir_ssa_def *global_id = get_global_ids(&b, 2); nir_ssa_def *src_offset = nir_load_push_constant(&b, 3, 32, nir_imm_int(&b, 0), .range = 24); nir_ssa_def *dst_offset = nir_load_push_constant(&b, 3, 32, nir_imm_int(&b, 12), .range = 24); @@ -943,13 +924,7 @@ build_nir_cleari_compute_shader(struct radv_device *dev, bool is_3d, int samples output_img->data.descriptor_set = 0; output_img->data.binding = 0; - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); + nir_ssa_def *global_id = get_global_ids(&b, 2); nir_ssa_def *clear_val = nir_load_push_constant(&b, 4, 32, nir_imm_int(&b, 0), .range = 20); nir_ssa_def *layer = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 16), .range = 20); @@ -1107,13 +1082,7 @@ build_nir_cleari_r32g32b32_compute_shader(struct radv_device *dev) output_img->data.descriptor_set = 0; output_img->data.binding = 0; - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); + nir_ssa_def *global_id = get_global_ids(&b, 2); nir_ssa_def *clear_val = nir_load_push_constant(&b, 3, 32, nir_imm_int(&b, 0), .range = 16); nir_ssa_def *stride = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 12), .range = 16); diff --git a/src/amd/vulkan/radv_meta_clear.c b/src/amd/vulkan/radv_meta_clear.c index 76f6bf46199..0974733fae2 100644 --- a/src/amd/vulkan/radv_meta_clear.c +++ b/src/amd/vulkan/radv_meta_clear.c @@ -1057,13 +1057,7 @@ build_clear_htile_mask_shader() b.shader->info.workgroup_size[1] = 1; b.shader->info.workgroup_size[2] = 1; - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); + nir_ssa_def *global_id = get_global_ids(&b, 1); nir_ssa_def *offset = nir_imul(&b, global_id, nir_imm_int(&b, 16)); offset = nir_channel(&b, offset, 0); @@ -1168,13 +1162,7 @@ build_clear_dcc_comp_to_single_shader(bool is_msaa) b.shader->info.workgroup_size[1] = 8; b.shader->info.workgroup_size[2] = 1; - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); - nir_ssa_def *layer_id = nir_channel(&b, wg_id, 2); + nir_ssa_def *global_id = get_global_ids(&b, 3); /* Load the dimensions in pixels of a block that gets compressed to one DCC byte. */ nir_ssa_def *dcc_block_size = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 8); @@ -1184,7 +1172,7 @@ build_clear_dcc_comp_to_single_shader(bool is_msaa) coord = nir_imul(&b, coord, dcc_block_size); coord = nir_vec4(&b, nir_channel(&b, coord, 0), nir_channel(&b, coord, 1), - layer_id, + nir_channel(&b, global_id, 2), nir_ssa_undef(&b, 1, 32)); nir_variable *output_img = nir_variable_create(b.shader, nir_var_uniform, img_type, "out_img"); diff --git a/src/amd/vulkan/radv_meta_copy_vrs_htile.c b/src/amd/vulkan/radv_meta_copy_vrs_htile.c index 27a6faba537..65d683e98ea 100644 --- a/src/amd/vulkan/radv_meta_copy_vrs_htile.c +++ b/src/amd/vulkan/radv_meta_copy_vrs_htile.c @@ -49,18 +49,11 @@ build_copy_vrs_htile_shader(struct radv_device *device, struct radeon_surf *surf b.shader->info.workgroup_size[1] = 8; b.shader->info.workgroup_size[2] = 1; - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - /* Get coordinates. */ - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); - nir_ssa_def *coord = nir_channels(&b, global_id, 0x3); + nir_ssa_def *global_id = get_global_ids(&b, 2); /* Multiply the coordinates by the HTILE block size. */ - coord = nir_imul(&b, coord, nir_imm_ivec2(&b, 8, 8)); + nir_ssa_def *coord = nir_imul(&b, global_id, nir_imm_ivec2(&b, 8, 8)); /* Load constants. */ nir_ssa_def *constants = nir_load_push_constant(&b, 3, 32, nir_imm_int(&b, 0), .range = 12); @@ -89,7 +82,7 @@ build_copy_vrs_htile_shader(struct radv_device *device, struct radeon_surf *surf tex->sampler_dim = GLSL_SAMPLER_DIM_2D; tex->op = nir_texop_txf; tex->src[0].src_type = nir_tex_src_coord; - tex->src[0].src = nir_src_for_ssa(nir_channels(&b, global_id, 0x3)); + tex->src[0].src = nir_src_for_ssa(global_id); tex->src[1].src_type = nir_tex_src_lod; tex->src[1].src = nir_src_for_ssa(nir_imm_int(&b, 0)); tex->src[2].src_type = nir_tex_src_texture_deref; diff --git a/src/amd/vulkan/radv_meta_dcc_retile.c b/src/amd/vulkan/radv_meta_dcc_retile.c index 943b41ce63f..1240015e8e0 100644 --- a/src/amd/vulkan/radv_meta_dcc_retile.c +++ b/src/amd/vulkan/radv_meta_dcc_retile.c @@ -27,22 +27,6 @@ #include "radv_meta.h" #include "radv_private.h" -static nir_ssa_def * -get_global_ids(nir_builder *b, unsigned num_components) -{ - unsigned mask = BITFIELD_MASK(num_components); - - nir_ssa_def *local_ids = nir_channels(b, nir_load_local_invocation_id(b), mask); - nir_ssa_def *block_ids = nir_channels(b, nir_load_workgroup_id(b, 32), mask); - nir_ssa_def *block_size = nir_channels( - b, - nir_imm_ivec4(b, b->shader->info.workgroup_size[0], b->shader->info.workgroup_size[1], - b->shader->info.workgroup_size[2], 0), - mask); - - return nir_iadd(b, nir_imul(b, block_ids, block_size), local_ids); -} - static nir_shader * build_dcc_retile_compute_shader(struct radv_device *dev, struct radeon_surf *surf) { diff --git a/src/amd/vulkan/radv_meta_fast_clear.c b/src/amd/vulkan/radv_meta_fast_clear.c index 0293a7e63e5..8b0673d7a32 100644 --- a/src/amd/vulkan/radv_meta_fast_clear.c +++ b/src/amd/vulkan/radv_meta_fast_clear.c @@ -54,16 +54,14 @@ build_dcc_decompress_compute_shader(struct radv_device *dev) output_img->data.descriptor_set = 0; output_img->data.binding = 1; - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); + nir_ssa_def *global_id = get_global_ids(&b, 2); + nir_ssa_def *img_coord = nir_vec4(&b, nir_channel(&b, global_id, 0), + nir_channel(&b, global_id, 1), + nir_ssa_undef(&b, 1, 32), + nir_ssa_undef(&b, 1, 32)); nir_ssa_def *data = nir_image_deref_load( - &b, 4, 32, &nir_build_deref_var(&b, input_img)->dest.ssa, global_id, nir_ssa_undef(&b, 1, 32), + &b, 4, 32, &nir_build_deref_var(&b, input_img)->dest.ssa, img_coord, nir_ssa_undef(&b, 1, 32), nir_imm_int(&b, 0), .image_dim = GLSL_SAMPLER_DIM_2D); /* We need a NIR_SCOPE_DEVICE memory_scope because ACO will avoid @@ -73,7 +71,7 @@ build_dcc_decompress_compute_shader(struct radv_device *dev) nir_scoped_barrier(&b, .execution_scope = NIR_SCOPE_WORKGROUP, .memory_scope = NIR_SCOPE_DEVICE, .memory_semantics = NIR_MEMORY_ACQ_REL, .memory_modes = nir_var_mem_ssbo); - nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->dest.ssa, global_id, + nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->dest.ssa, img_coord, nir_ssa_undef(&b, 1, 32), data, nir_imm_int(&b, 0), .image_dim = GLSL_SAMPLER_DIM_2D); return b.shader; diff --git a/src/amd/vulkan/radv_meta_fmask_expand.c b/src/amd/vulkan/radv_meta_fmask_expand.c index 0ac11d547f3..b4c1276b60e 100644 --- a/src/amd/vulkan/radv_meta_fmask_expand.c +++ b/src/amd/vulkan/radv_meta_fmask_expand.c @@ -48,20 +48,10 @@ build_fmask_expand_compute_shader(struct radv_device *device, int samples) output_img->data.binding = 1; output_img->data.access = ACCESS_NON_READABLE; - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); - nir_ssa_def *layer_id = nir_channel(&b, wg_id, 2); - nir_ssa_def *input_img_deref = &nir_build_deref_var(&b, input_img)->dest.ssa; nir_ssa_def *output_img_deref = &nir_build_deref_var(&b, output_img)->dest.ssa; - nir_ssa_def *tex_coord = - nir_vec3(&b, nir_channel(&b, global_id, 0), nir_channel(&b, global_id, 1), layer_id); + nir_ssa_def *tex_coord = get_global_ids(&b, 3); nir_tex_instr *tex_instr[8]; for (uint32_t i = 0; i < samples; i++) { diff --git a/src/amd/vulkan/radv_meta_resolve_cs.c b/src/amd/vulkan/radv_meta_resolve_cs.c index 69ecb6cd208..190d6204b3c 100644 --- a/src/amd/vulkan/radv_meta_resolve_cs.c +++ b/src/amd/vulkan/radv_meta_resolve_cs.c @@ -78,28 +78,29 @@ build_resolve_compute_shader(struct radv_device *dev, bool is_integer, bool is_s nir_variable *output_img = nir_variable_create(b.shader, nir_var_uniform, img_type, "out_img"); output_img->data.descriptor_set = 0; output_img->data.binding = 1; - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); + nir_ssa_def *global_id = get_global_ids(&b, 2); nir_ssa_def *src_offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 16); nir_ssa_def *dst_offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 8), .range = 16); - nir_ssa_def *img_coord = nir_channels(&b, nir_iadd(&b, global_id, src_offset), 0x3); + nir_ssa_def *src_coord = nir_iadd(&b, global_id, src_offset); + nir_ssa_def *dst_coord = nir_iadd(&b, global_id, dst_offset); + nir_variable *color = nir_local_variable_create(b.impl, glsl_vec4_type(), "color"); - radv_meta_build_resolve_shader_core(&b, is_integer, samples, input_img, color, img_coord); + radv_meta_build_resolve_shader_core(&b, is_integer, samples, input_img, color, src_coord); nir_ssa_def *outval = nir_load_var(&b, color); if (is_srgb) outval = radv_meta_build_resolve_srgb_conversion(&b, outval); - nir_ssa_def *coord = nir_iadd(&b, global_id, dst_offset); - nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->dest.ssa, coord, + nir_ssa_def *img_coord = nir_vec4(&b, nir_channel(&b, dst_coord, 0), + nir_channel(&b, dst_coord, 1), + nir_ssa_undef(&b, 1, 32), + nir_ssa_undef(&b, 1, 32)); + + nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->dest.ssa, img_coord, nir_ssa_undef(&b, 1, 32), outval, nir_imm_int(&b, 0), .image_dim = GLSL_SAMPLER_DIM_2D); return b.shader; @@ -149,17 +150,8 @@ build_depth_stencil_resolve_compute_shader(struct radv_device *dev, int samples, nir_variable *output_img = nir_variable_create(b.shader, nir_var_uniform, img_type, "out_img"); output_img->data.descriptor_set = 0; output_img->data.binding = 1; - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); - nir_ssa_def *layer_id = nir_channel(&b, wg_id, 2); - - nir_ssa_def *img_coord = - nir_vec3(&b, nir_channel(&b, global_id, 0), nir_channel(&b, global_id, 1), layer_id); + nir_ssa_def *img_coord = get_global_ids(&b, 3); nir_ssa_def *input_img_deref = &nir_build_deref_var(&b, input_img)->dest.ssa; diff --git a/src/amd/vulkan/radv_query.c b/src/amd/vulkan/radv_query.c index af1d4a4a1e2..73d005111a7 100644 --- a/src/amd/vulkan/radv_query.c +++ b/src/amd/vulkan/radv_query.c @@ -149,13 +149,7 @@ build_occlusion_query_shader(struct radv_device *device) nir_ssa_def *dst_buf = radv_meta_load_descriptor(&b, 0, 0); nir_ssa_def *src_buf = radv_meta_load_descriptor(&b, 0, 1); - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); - global_id = nir_channel(&b, global_id, 0); // We only care about x here. + nir_ssa_def *global_id = get_global_ids(&b, 1); nir_ssa_def *input_stride = nir_imm_int(&b, db_count * 16); nir_ssa_def *input_base = nir_imul(&b, input_stride, global_id); @@ -290,13 +284,7 @@ build_pipeline_statistics_query_shader(struct radv_device *device) nir_ssa_def *dst_buf = radv_meta_load_descriptor(&b, 0, 0); nir_ssa_def *src_buf = radv_meta_load_descriptor(&b, 0, 1); - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); - global_id = nir_channel(&b, global_id, 0); // We only care about x here. + nir_ssa_def *global_id = get_global_ids(&b, 1); nir_ssa_def *input_stride = nir_imm_int(&b, pipelinestat_block_size * 2); nir_ssa_def *input_base = nir_imul(&b, input_stride, global_id); @@ -441,13 +429,7 @@ build_tfb_query_shader(struct radv_device *device) nir_ssa_def *src_buf = radv_meta_load_descriptor(&b, 0, 1); /* Compute global ID. */ - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); - global_id = nir_channel(&b, global_id, 0); // We only care about x here. + nir_ssa_def *global_id = get_global_ids(&b, 1); /* Compute src/dst strides. */ nir_ssa_def *input_stride = nir_imm_int(&b, 32); @@ -571,13 +553,7 @@ build_timestamp_query_shader(struct radv_device *device) nir_ssa_def *src_buf = radv_meta_load_descriptor(&b, 0, 1); /* Compute global ID. */ - nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b); - nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32); - nir_ssa_def *block_size = - nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1], - b.shader->info.workgroup_size[2], 0); - nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id); - global_id = nir_channel(&b, global_id, 0); // We only care about x here. + nir_ssa_def *global_id = get_global_ids(&b, 1); /* Compute src/dst strides. */ nir_ssa_def *input_stride = nir_imm_int(&b, 8);