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

radv: use get_global_ids() to compute coordinates in meta shaders

Browse source 
This was duplicated everywhere.

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12561>
This commit is contained in:
Samuel Pitoiset 2021-08-26 09:02:07 +02:00 • committed by Marge Bot
parent 80df2110b1
commit f741c04ed1
11 changed files with 70 additions and 174 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

16
src/amd/vulkan/radv_meta.c
View file

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

2
src/amd/vulkan/radv_meta.h
View file

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

16
src/amd/vulkan/radv_meta_buffer.c
View file

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

71
src/amd/vulkan/radv_meta_bufimage.c
View file

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

18
src/amd/vulkan/radv_meta_clear.c
View file

@ -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");

13
src/amd/vulkan/radv_meta_copy_vrs_htile.c
View file

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

16
src/amd/vulkan/radv_meta_dcc_retile.c
View file

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

16
src/amd/vulkan/radv_meta_fast_clear.c
View file

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

12
src/amd/vulkan/radv_meta_fmask_expand.c
View file

@ -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++) {

32
src/amd/vulkan/radv_meta_resolve_cs.c
View file

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

32
src/amd/vulkan/radv_query.c
View file

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