Mesa (main): radv: use get_global_ids() to compute coordinates in meta shaders
GitLab Mirror
gitlab-mirror at kemper.freedesktop.org
Mon Oct 11 10:29:57 UTC 2021
Module: Mesa
Branch: main
Commit: f741c04ed10a5ebaebf786ba5cd8b75581fc08a1
URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=f741c04ed10a5ebaebf786ba5cd8b75581fc08a1
Author: Samuel Pitoiset <samuel.pitoiset at gmail.com>
Date: Thu Aug 26 09:02:07 2021 +0200
radv: use get_global_ids() to compute coordinates in meta shaders
This was duplicated everywhere.
Signed-off-by: Samuel Pitoiset <samuel.pitoiset at gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas at basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12561>
---
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);
More information about the mesa-commit
mailing list