fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2025-12-20 07:20:10 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 2

radv: Rename get_global_ids to radv_meta_nir_get_global_ids.

Browse source 
Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/33494>
This commit is contained in:
Timur Kristóf 2025-02-11 18:59:44 +01:00 • committed by Marge Bot
parent 09db738c9a
commit 6b57cc2fbf
4 changed files with 27 additions and 27 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

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

@ -103,7 +103,7 @@ radv_meta_nir_load_descriptor(nir_builder *b, unsigned desc_set, unsigned bindin
}
nir_def *
get_global_ids(nir_builder *b, unsigned num_components)
radv_meta_nir_get_global_ids(nir_builder *b, unsigned num_components)
{
unsigned mask = BITFIELD_MASK(num_components);
@ -501,7 +501,7 @@ radv_meta_nir_build_itob_compute_shader(struct radv_device *dev, bool is_3d)
output_img->data.descriptor_set = 0;
output_img->data.binding = 1;
nir_def *global_id = get_global_ids(&b, is_3d ? 3 : 2);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, is_3d ? 3 : 2);
nir_def *offset = nir_load_push_constant(&b, is_3d ? 3 : 2, 32, nir_imm_int(&b, 0), .range = is_3d ? 12 : 8);
nir_def *stride = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 12), .range = 16);
@ -541,7 +541,7 @@ radv_meta_nir_build_btoi_compute_shader(struct radv_device *dev, bool is_3d)
output_img->data.descriptor_set = 0;
output_img->data.binding = 1;
nir_def *global_id = get_global_ids(&b, is_3d ? 3 : 2);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, is_3d ? 3 : 2);
nir_def *offset = nir_load_push_constant(&b, is_3d ? 3 : 2, 32, nir_imm_int(&b, 0), .range = is_3d ? 12 : 8);
nir_def *stride = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 12), .range = 16);
@ -581,7 +581,7 @@ radv_meta_nir_build_btoi_r32g32b32_compute_shader(struct radv_device *dev)
output_img->data.descriptor_set = 0;
output_img->data.binding = 1;
nir_def *global_id = get_global_ids(&b, 2);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 2);
nir_def *offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 8);
nir_def *pitch = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 8), .range = 12);
@ -636,7 +636,7 @@ radv_meta_nir_build_itoi_compute_shader(struct radv_device *dev, bool src_3d, bo
output_img->data.descriptor_set = 0;
output_img->data.binding = 1;
nir_def *global_id = get_global_ids(&b, (src_3d || dst_3d) ? 3 : 2);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, (src_3d || dst_3d) ? 3 : 2);
nir_def *src_offset = nir_load_push_constant(&b, src_3d ? 3 : 2, 32, nir_imm_int(&b, 0), .range = src_3d ? 12 : 8);
nir_def *dst_offset = nir_load_push_constant(&b, dst_3d ? 3 : 2, 32, nir_imm_int(&b, 12), .range = dst_3d ? 24 : 20);
@ -682,7 +682,7 @@ radv_meta_nir_build_itoi_r32g32b32_compute_shader(struct radv_device *dev)
output_img->data.descriptor_set = 0;
output_img->data.binding = 1;
nir_def *global_id = get_global_ids(&b, 2);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 2);
nir_def *src_offset = nir_load_push_constant(&b, 3, 32, nir_imm_int(&b, 0), .range = 12);
nir_def *dst_offset = nir_load_push_constant(&b, 3, 32, nir_imm_int(&b, 12), .range = 24);
@ -733,7 +733,7 @@ radv_meta_nir_build_cleari_compute_shader(struct radv_device *dev, bool is_3d, i
output_img->data.descriptor_set = 0;
output_img->data.binding = 0;
nir_def *global_id = get_global_ids(&b, 2);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 2);
nir_def *clear_val = nir_load_push_constant(&b, 4, 32, nir_imm_int(&b, 0), .range = 16);
nir_def *layer = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 16), .range = 20);
@ -766,7 +766,7 @@ radv_meta_nir_build_cleari_r32g32b32_compute_shader(struct radv_device *dev)
output_img->data.descriptor_set = 0;
output_img->data.binding = 0;
nir_def *global_id = get_global_ids(&b, 2);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 2);
nir_def *clear_val = nir_load_push_constant(&b, 3, 32, nir_imm_int(&b, 0), .range = 12);
nir_def *stride = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 12), .range = 16);
@ -876,7 +876,7 @@ radv_meta_nir_build_clear_htile_mask_shader(struct radv_device *dev)
nir_builder b = radv_meta_nir_init_shader(dev, MESA_SHADER_COMPUTE, "meta_clear_htile_mask");
b.shader->info.workgroup_size[0] = 64;
nir_def *global_id = get_global_ids(&b, 1);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 1);
nir_def *offset = nir_imul_imm(&b, global_id, 16);
offset = nir_channel(&b, offset, 0);
@ -911,7 +911,7 @@ radv_meta_nir_build_clear_dcc_comp_to_single_shader(struct radv_device *dev, boo
b.shader->info.workgroup_size[0] = 8;
b.shader->info.workgroup_size[1] = 8;
nir_def *global_id = get_global_ids(&b, 3);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 3);
/* Load the dimensions in pixels of a block that gets compressed to one DCC byte. */
nir_def *dcc_block_size = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 8);
@ -949,7 +949,7 @@ radv_meta_nir_build_copy_vrs_htile_shader(struct radv_device *device, struct rad
b.shader->info.workgroup_size[1] = 8;
/* Get coordinates. */
nir_def *global_id = get_global_ids(&b, 2);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 2);
nir_def *addr = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 8);
nir_def *htile_va = nir_pack_64_2x32(&b, nir_channels(&b, addr, 0x3));
@ -1050,7 +1050,7 @@ radv_meta_nir_build_dcc_retile_compute_shader(struct radv_device *dev, struct ra
nir_def *input_dcc_ref = &nir_build_deref_var(&b, input_dcc)->def;
nir_def *output_dcc_ref = &nir_build_deref_var(&b, output_dcc)->def;
nir_def *coord = get_global_ids(&b, 2);
nir_def *coord = radv_meta_nir_get_global_ids(&b, 2);
nir_def *zero = nir_imm_int(&b, 0);
coord =
nir_imul(&b, coord, nir_imm_ivec2(&b, surf->u.gfx9.color.dcc_block_width, surf->u.gfx9.color.dcc_block_height));
@ -1129,7 +1129,7 @@ radv_meta_nir_build_dcc_decompress_compute_shader(struct radv_device *dev)
output_img->data.descriptor_set = 0;
output_img->data.binding = 1;
nir_def *global_id = get_global_ids(&b, 2);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 2);
nir_def *img_coord = nir_vec4(&b, nir_channel(&b, global_id, 0), nir_channel(&b, global_id, 1), nir_undef(&b, 1, 32),
nir_undef(&b, 1, 32));
@ -1245,7 +1245,7 @@ radv_meta_nir_build_fmask_expand_compute_shader(struct radv_device *device, int
nir_deref_instr *input_img_deref = nir_build_deref_var(&b, input_img);
nir_def *output_img_deref = &nir_build_deref_var(&b, output_img)->def;
nir_def *tex_coord = get_global_ids(&b, 3);
nir_def *tex_coord = radv_meta_nir_get_global_ids(&b, 3);
nir_def *tex_vals[8];
for (uint32_t i = 0; i < samples; i++) {
@ -1293,7 +1293,7 @@ radv_meta_nir_build_resolve_compute_shader(struct radv_device *dev, bool is_inte
output_img->data.descriptor_set = 0;
output_img->data.binding = 1;
nir_def *global_id = get_global_ids(&b, 2);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 2);
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);
@ -1357,7 +1357,7 @@ radv_meta_nir_build_depth_stencil_resolve_compute_shader(struct radv_device *dev
output_img->data.descriptor_set = 0;
output_img->data.binding = 1;
nir_def *global_id = get_global_ids(&b, 3);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 3);
nir_def *offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 8);

2
src/amd/vulkan/nir/radv_meta_nir.h
View file

@ -26,7 +26,7 @@ nir_builder PRINTFLIKE(3, 4)
nir_shader *radv_meta_nir_build_vs_generate_vertices(struct radv_device *dev);
nir_shader *radv_meta_nir_build_fs_noop(struct radv_device *dev);
nir_def *get_global_ids(nir_builder *b, unsigned num_components);
nir_def *radv_meta_nir_get_global_ids(nir_builder *b, unsigned num_components);
void radv_meta_nir_break_on_count(nir_builder *b, nir_variable *var, nir_def *count);

8
src/amd/vulkan/radv_dgc.c
View file

@ -929,7 +929,7 @@ build_dgc_buffer_tail(nir_builder *b, nir_def *cmd_buf_offset, nir_def *cmd_buf_
const struct radv_physical_device *pdev = radv_device_physical(device);
nir_def *is_compute_queue = nir_ior_imm(b, nir_ieq_imm(b, load_param8(b, queue_family), RADV_QUEUE_COMPUTE), is_ace);
nir_def *global_id = get_global_ids(b, 1);
nir_def *global_id = radv_meta_nir_get_global_ids(b, 1);
nir_push_if(b, nir_ieq_imm(b, global_id, 0));
{
@ -1018,7 +1018,7 @@ build_dgc_buffer_trailer(nir_builder *b, nir_def *cmd_buf_offset, unsigned trail
{
const struct radv_physical_device *pdev = radv_device_physical(device);
nir_def *global_id = get_global_ids(b, 1);
nir_def *global_id = radv_meta_nir_get_global_ids(b, 1);
nir_push_if(b, nir_ieq_imm(b, global_id, 0));
{
@ -1071,7 +1071,7 @@ build_dgc_buffer_preamble(nir_builder *b, nir_def *cmd_buf_preamble_offset, nir_
{
const struct radv_physical_device *pdev = radv_device_physical(device);
nir_def *global_id = get_global_ids(b, 1);
nir_def *global_id = radv_meta_nir_get_global_ids(b, 1);
nir_def *use_preamble = nir_ine_imm(b, load_param8(b, use_preamble), 0);
nir_push_if(b, nir_iand(b, nir_ieq_imm(b, global_id, 0), use_preamble));
@ -2474,7 +2474,7 @@ build_dgc_prepare_shader(struct radv_device *dev, struct radv_indirect_command_l
nir_builder b = radv_meta_nir_init_shader(dev, MESA_SHADER_COMPUTE, "meta_dgc_prepare");
b.shader->info.workgroup_size[0] = 64;
nir_def *global_id = get_global_ids(&b, 1);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 1);
nir_def *sequence_id = global_id;

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

@ -173,7 +173,7 @@ build_occlusion_query_shader(struct radv_device *device)
nir_def *flags = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 16), .range = 20);
nir_def *global_id = get_global_ids(&b, 1);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 1);
nir_def *input_stride = nir_imm_int(&b, db_count * 16);
nir_def *input_base = nir_imul(&b, input_stride, global_id);
@ -464,7 +464,7 @@ build_pipeline_statistics_query_shader(struct radv_device *device)
nir_def *avail_offset = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 28), .range = 32);
nir_def *uses_emulated_queries = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 32), .range = 36);
nir_def *global_id = get_global_ids(&b, 1);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 1);
nir_def *input_stride =
nir_bcsel(&b, nir_ine_imm(&b, uses_emulated_queries, 0), nir_imm_int(&b, pipelinestat_block_size * 2 + 8 * 2),
@ -860,7 +860,7 @@ build_tfb_query_shader(struct radv_device *device)
nir_def *dst_va = nir_pack_64_2x32(&b, nir_channels(&b, addrs, 0xc));
/* Compute global ID. */
nir_def *global_id = get_global_ids(&b, 1);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 1);
/* Compute src/dst strides. */
nir_def *input_stride = nir_imm_int(&b, 32);
@ -1123,7 +1123,7 @@ build_timestamp_query_shader(struct radv_device *device)
nir_def *dst_va = nir_pack_64_2x32(&b, nir_channels(&b, addrs, 0xc));
/* Compute global ID. */
nir_def *global_id = get_global_ids(&b, 1);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 1);
/* Compute src/dst strides. */
nir_def *input_stride = nir_imm_int(&b, 8);
@ -1267,7 +1267,7 @@ build_pg_query_shader(struct radv_device *device)
nir_def *dst_va = nir_pack_64_2x32(&b, nir_channels(&b, addrs, 0xc));
/* Compute global ID. */
nir_def *global_id = get_global_ids(&b, 1);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 1);
/* Determine if the query pool uses emulated queries for NGG. */
nir_def *uses_emulated_queries = nir_i2b(&b, nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 32), .range = 36));
@ -1548,7 +1548,7 @@ build_ms_prim_gen_query_shader(struct radv_device *device)
nir_def *dst_va = nir_pack_64_2x32(&b, nir_channels(&b, addrs, 0xc));
/* Compute global ID. */
nir_def *global_id = get_global_ids(&b, 1);
nir_def *global_id = radv_meta_nir_get_global_ids(&b, 1);
/* Compute src/dst strides. */
nir_def *input_base = nir_imul_imm(&b, global_id, 16);