KonstantinSeurer
/
mesa
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

compiler: Rename local_size to workgroup_size 

Acked-by: Emma Anholt <emma@anholt.net>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Acked-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11190>
This commit is contained in:
Caio Marcelo de Oliveira Filho 2021-05-26 23:53:32 -07:00 committed by Marge Bot
parent 4b9e52e818
commit 430d2206da
49 changed files with 286 additions and 286 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
src/amd/compiler/aco_instruction_selection_setup.cpp
View File

@ -1078,9 +1078,9 @@ setup_isel_context(Program* program,
program->workgroup_size = program->wave_size;
} else if (program->stage == compute_cs) {
/* CS sets the workgroup size explicitly */
program->workgroup_size = shaders[0]->info.cs.local_size[0] *
shaders[0]->info.cs.local_size[1] *
shaders[0]->info.cs.local_size[2];
program->workgroup_size = shaders[0]->info.cs.workgroup_size[0] *
shaders[0]->info.cs.workgroup_size[1] *
shaders[0]->info.cs.workgroup_size[2];
} else if (program->stage.hw == HWStage::ES || program->stage == geometry_gs) {
/* Unmerged ESGS operate in workgroups if on-chip GS (LDS rings) are enabled on GFX7-8 (not implemented in Mesa) */
program->workgroup_size = program->wave_size;

20
src/amd/vulkan/radv_meta_buffer.c
View File

@ -8,15 +8,15 @@ static nir_shader *
build_buffer_fill_shader(struct radv_device *dev)
{
nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "meta_buffer_fill");
b.shader->info.cs.local_size[0] = 64;
b.shader->info.cs.local_size[1] = 1;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 64;
b.shader->info.cs.workgroup_size[1] = 1;
b.shader->info.cs.workgroup_size[2] = 1;
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.workgroup_size[2], 0);
nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id);
@ -38,15 +38,15 @@ static nir_shader *
build_buffer_copy_shader(struct radv_device *dev)
{
nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "meta_buffer_copy");
b.shader->info.cs.local_size[0] = 64;
b.shader->info.cs.local_size[1] = 1;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 64;
b.shader->info.cs.workgroup_size[1] = 1;
b.shader->info.cs.workgroup_size[2] = 1;
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.workgroup_size[2], 0);
nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id);

70
src/amd/vulkan/radv_meta_bufimage.c
View File

@ -40,9 +40,9 @@ build_nir_itob_compute_shader(struct radv_device *dev, bool is_3d)
const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_BUF, false, GLSL_TYPE_FLOAT);
nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL,
is_3d ? "meta_itob_cs_3d" : "meta_itob_cs");
b.shader->info.cs.local_size[0] = 8;
b.shader->info.cs.local_size[1] = 8;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 8;
b.shader->info.cs.workgroup_size[1] = 8;
b.shader->info.cs.workgroup_size[2] = 1;
nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, sampler_type, "s_tex");
input_img->data.descriptor_set = 0;
input_img->data.binding = 0;
@ -54,8 +54,8 @@ build_nir_itob_compute_shader(struct radv_device *dev, bool is_3d)
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.workgroup_size[2], 0);
nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id);
@ -227,9 +227,9 @@ build_nir_btoi_compute_shader(struct radv_device *dev, bool is_3d)
const struct glsl_type *img_type = glsl_image_type(dim, false, GLSL_TYPE_FLOAT);
nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL,
is_3d ? "meta_btoi_cs_3d" : "meta_btoi_cs");
b.shader->info.cs.local_size[0] = 8;
b.shader->info.cs.local_size[1] = 8;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 8;
b.shader->info.cs.workgroup_size[1] = 8;
b.shader->info.cs.workgroup_size[2] = 1;
nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, buf_type, "s_tex");
input_img->data.descriptor_set = 0;
input_img->data.binding = 0;
@ -241,8 +241,8 @@ build_nir_btoi_compute_shader(struct radv_device *dev, bool is_3d)
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.workgroup_size[2], 0);
nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id);
@ -409,9 +409,9 @@ build_nir_btoi_r32g32b32_compute_shader(struct radv_device *dev)
const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_BUF, false, GLSL_TYPE_FLOAT);
nir_builder b =
nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "meta_btoi_r32g32b32_cs");
b.shader->info.cs.local_size[0] = 8;
b.shader->info.cs.local_size[1] = 8;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 8;
b.shader->info.cs.workgroup_size[1] = 8;
b.shader->info.cs.workgroup_size[2] = 1;
nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, buf_type, "s_tex");
input_img->data.descriptor_set = 0;
input_img->data.binding = 0;
@ -423,8 +423,8 @@ build_nir_btoi_r32g32b32_compute_shader(struct radv_device *dev)
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.workgroup_size[2], 0);
nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id);
@ -571,9 +571,9 @@ build_nir_itoi_compute_shader(struct radv_device *dev, bool is_3d, int samples)
const struct glsl_type *img_type = glsl_image_type(dim, false, GLSL_TYPE_FLOAT);
nir_builder b = nir_builder_init_simple_shader(
MESA_SHADER_COMPUTE, NULL, is_3d ? "meta_itoi_cs_3d-%d" : "meta_itoi_cs-%d", samples);
b.shader->info.cs.local_size[0] = 8;
b.shader->info.cs.local_size[1] = 8;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 8;
b.shader->info.cs.workgroup_size[1] = 8;
b.shader->info.cs.workgroup_size[2] = 1;
nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, buf_type, "s_tex");
input_img->data.descriptor_set = 0;
input_img->data.binding = 0;
@ -585,8 +585,8 @@ build_nir_itoi_compute_shader(struct radv_device *dev, bool is_3d, int samples)
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.workgroup_size[2], 0);
nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id);
@ -772,9 +772,9 @@ build_nir_itoi_r32g32b32_compute_shader(struct radv_device *dev)
const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_BUF, false, GLSL_TYPE_FLOAT);
nir_builder b =
nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "meta_itoi_r32g32b32_cs");
b.shader->info.cs.local_size[0] = 8;
b.shader->info.cs.local_size[1] = 8;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 8;
b.shader->info.cs.workgroup_size[1] = 8;
b.shader->info.cs.workgroup_size[2] = 1;
nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, type, "input_img");
input_img->data.descriptor_set = 0;
input_img->data.binding = 0;
@ -787,8 +787,8 @@ build_nir_itoi_r32g32b32_compute_shader(struct radv_device *dev)
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.workgroup_size[2], 0);
nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id);
@ -942,9 +942,9 @@ build_nir_cleari_compute_shader(struct radv_device *dev, bool is_3d, int samples
const struct glsl_type *img_type = glsl_image_type(dim, false, GLSL_TYPE_FLOAT);
nir_builder b = nir_builder_init_simple_shader(
MESA_SHADER_COMPUTE, NULL, is_3d ? "meta_cleari_cs_3d-%d" : "meta_cleari_cs-%d", samples);
b.shader->info.cs.local_size[0] = 8;
b.shader->info.cs.local_size[1] = 8;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 8;
b.shader->info.cs.workgroup_size[1] = 8;
b.shader->info.cs.workgroup_size[2] = 1;
nir_variable *output_img = nir_variable_create(b.shader, nir_var_uniform, img_type, "out_img");
output_img->data.descriptor_set = 0;
@ -953,8 +953,8 @@ build_nir_cleari_compute_shader(struct radv_device *dev, bool is_3d, int samples
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.workgroup_size[2], 0);
nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id);
@ -1108,9 +1108,9 @@ build_nir_cleari_r32g32b32_compute_shader(struct radv_device *dev)
const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_BUF, false, GLSL_TYPE_FLOAT);
nir_builder b =
nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "meta_cleari_r32g32b32_cs");
b.shader->info.cs.local_size[0] = 8;
b.shader->info.cs.local_size[1] = 8;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 8;
b.shader->info.cs.workgroup_size[1] = 8;
b.shader->info.cs.workgroup_size[2] = 1;
nir_variable *output_img = nir_variable_create(b.shader, nir_var_uniform, img_type, "out_img");
output_img->data.descriptor_set = 0;
@ -1119,8 +1119,8 @@ build_nir_cleari_r32g32b32_compute_shader(struct radv_device *dev)
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.workgroup_size[2], 0);
nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id);

10
src/amd/vulkan/radv_meta_clear.c
View File

@ -1002,15 +1002,15 @@ build_clear_htile_mask_shader()
{
nir_builder b =
nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "meta_clear_htile_mask");
b.shader->info.cs.local_size[0] = 64;
b.shader->info.cs.local_size[1] = 1;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 64;
b.shader->info.cs.workgroup_size[1] = 1;
b.shader->info.cs.workgroup_size[2] = 1;
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.workgroup_size[2], 0);
nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id);

10
src/amd/vulkan/radv_meta_copy_vrs_htile.c
View File

@ -45,15 +45,15 @@ static nir_shader *
build_copy_vrs_htile_shader(struct radv_device *device, struct radeon_surf *surf)
{
nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "meta_copy_vrs_htile");
b.shader->info.cs.local_size[0] = 8;
b.shader->info.cs.local_size[1] = 8;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 8;
b.shader->info.cs.workgroup_size[1] = 8;
b.shader->info.cs.workgroup_size[2] = 1;
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.workgroup_size[2], 0);
/* Get coordinates. */
nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id);

10
src/amd/vulkan/radv_meta_dcc_retile.c
View File

@ -36,8 +36,8 @@ get_global_ids(nir_builder *b, unsigned num_components)
nir_ssa_def *block_ids = nir_channels(b, nir_load_work_group_id(b, 32), mask);
nir_ssa_def *block_size = nir_channels(
b,
nir_imm_ivec4(b, b->shader->info.cs.local_size[0], b->shader->info.cs.local_size[1],
b->shader->info.cs.local_size[2], 0),
nir_imm_ivec4(b, b->shader->info.cs.workgroup_size[0], b->shader->info.cs.workgroup_size[1],
b->shader->info.cs.workgroup_size[2], 0),
mask);
return nir_iadd(b, nir_imul(b, block_ids, block_size), local_ids);
@ -49,9 +49,9 @@ build_dcc_retile_compute_shader(struct radv_device *dev, struct radeon_surf *sur
const struct glsl_type *buf_type = glsl_image_type(GLSL_SAMPLER_DIM_BUF, false, GLSL_TYPE_UINT);
nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "dcc_retile_compute");
b.shader->info.cs.local_size[0] = 8;
b.shader->info.cs.local_size[1] = 8;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 8;
b.shader->info.cs.workgroup_size[1] = 8;
b.shader->info.cs.workgroup_size[2] = 1;
nir_ssa_def *src_dcc_size = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 8);
nir_ssa_def *src_dcc_pitch = nir_channels(&b, src_dcc_size, 1);

10
src/amd/vulkan/radv_meta_fast_clear.c
View File

@ -37,9 +37,9 @@ build_dcc_decompress_compute_shader(struct radv_device *dev)
nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "dcc_decompress_compute");
/* We need at least 16/16/1 to cover an entire DCC block in a single workgroup. */
b.shader->info.cs.local_size[0] = 16;
b.shader->info.cs.local_size[1] = 16;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 16;
b.shader->info.cs.workgroup_size[1] = 16;
b.shader->info.cs.workgroup_size[2] = 1;
nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, img_type, "in_img");
input_img->data.descriptor_set = 0;
input_img->data.binding = 0;
@ -51,8 +51,8 @@ build_dcc_decompress_compute_shader(struct radv_device *dev)
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.workgroup_size[2], 0);
nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id);

10
src/amd/vulkan/radv_meta_fmask_expand.c
View File

@ -35,9 +35,9 @@ build_fmask_expand_compute_shader(struct radv_device *device, int samples)
nir_builder b =
nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "meta_fmask_expand_cs-%d", samples);
b.shader->info.cs.local_size[0] = 8;
b.shader->info.cs.local_size[1] = 8;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 8;
b.shader->info.cs.workgroup_size[1] = 8;
b.shader->info.cs.workgroup_size[2] = 1;
nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, type, "s_tex");
input_img->data.descriptor_set = 0;
@ -51,8 +51,8 @@ build_fmask_expand_compute_shader(struct radv_device *device, int samples)
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.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);

20
src/amd/vulkan/radv_meta_resolve_cs.c
View File

@ -67,9 +67,9 @@ build_resolve_compute_shader(struct radv_device *dev, bool is_integer, bool is_s
nir_builder b =
nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "meta_resolve_cs-%d-%s", samples,
is_integer ? "int" : (is_srgb ? "srgb" : "float"));
b.shader->info.cs.local_size[0] = 8;
b.shader->info.cs.local_size[1] = 8;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 8;
b.shader->info.cs.workgroup_size[1] = 8;
b.shader->info.cs.workgroup_size[2] = 1;
nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, sampler_type, "s_tex");
input_img->data.descriptor_set = 0;
@ -81,8 +81,8 @@ build_resolve_compute_shader(struct radv_device *dev, bool is_integer, bool is_s
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.workgroup_size[2], 0);
nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id);
@ -137,9 +137,9 @@ build_depth_stencil_resolve_compute_shader(struct radv_device *dev, int samples,
nir_builder b = nir_builder_init_simple_shader(
MESA_SHADER_COMPUTE, NULL, "meta_resolve_cs_%s-%s-%d",
index == DEPTH_RESOLVE ? "depth" : "stencil", get_resolve_mode_str(resolve_mode), samples);
b.shader->info.cs.local_size[0] = 8;
b.shader->info.cs.local_size[1] = 8;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 8;
b.shader->info.cs.workgroup_size[1] = 8;
b.shader->info.cs.workgroup_size[2] = 1;
nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, sampler_type, "s_tex");
input_img->data.descriptor_set = 0;
@ -151,8 +151,8 @@ build_depth_stencil_resolve_compute_shader(struct radv_device *dev, int samples,
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.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);

2
src/amd/vulkan/radv_nir_to_llvm.c
View File

@ -2935,7 +2935,7 @@ radv_nir_get_max_workgroup_size(enum chip_class chip_class, gl_shader_stage stag
const unsigned backup_sizes[] = {chip_class >= GFX9 ? 128 : 64, 1, 1};
unsigned sizes[3];
for (unsigned i = 0; i < 3; i++)
sizes[i] = nir ? nir->info.cs.local_size[i] : backup_sizes[i];
sizes[i] = nir ? nir->info.cs.workgroup_size[i] : backup_sizes[i];
return radv_get_max_workgroup_size(chip_class, stage, sizes);
}

6
src/amd/vulkan/radv_pipeline.c
View File

@ -2791,9 +2791,9 @@ radv_fill_shader_keys(struct radv_device *device, struct radv_shader_variant_key
if (!subgroup_size)
subgroup_size = device->physical_device->cs_wave_size;
unsigned local_size = nir[MESA_SHADER_COMPUTE]->info.cs.local_size[0] *
nir[MESA_SHADER_COMPUTE]->info.cs.local_size[1] *
nir[MESA_SHADER_COMPUTE]->info.cs.local_size[2];
unsigned local_size = nir[MESA_SHADER_COMPUTE]->info.cs.workgroup_size[0] *
nir[MESA_SHADER_COMPUTE]->info.cs.workgroup_size[1] *
nir[MESA_SHADER_COMPUTE]->info.cs.workgroup_size[2];
/* Games don't always request full subgroups when they should,
* which can cause bugs if cswave32 is enabled.

40
src/amd/vulkan/radv_query.c
View File

@ -130,9 +130,9 @@ build_occlusion_query_shader(struct radv_device *device)
* }
*/
nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "occlusion_query");
b.shader->info.cs.local_size[0] = 64;
b.shader->info.cs.local_size[1] = 1;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 64;
b.shader->info.cs.workgroup_size[1] = 1;
b.shader->info.cs.workgroup_size[2] = 1;
nir_variable *result = nir_local_variable_create(b.impl, glsl_uint64_t_type(), "result");
nir_variable *outer_counter =
@ -151,8 +151,8 @@ build_occlusion_query_shader(struct radv_device *device)
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.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.
@ -275,9 +275,9 @@ build_pipeline_statistics_query_shader(struct radv_device *device)
*/
nir_builder b =
nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "pipeline_statistics_query");
b.shader->info.cs.local_size[0] = 64;
b.shader->info.cs.local_size[1] = 1;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 64;
b.shader->info.cs.workgroup_size[1] = 1;
b.shader->info.cs.workgroup_size[2] = 1;
nir_variable *output_offset =
nir_local_variable_create(b.impl, glsl_int_type(), "output_offset");
@ -292,8 +292,8 @@ build_pipeline_statistics_query_shader(struct radv_device *device)
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.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.
@ -421,9 +421,9 @@ build_tfb_query_shader(struct radv_device *device)
* }
*/
nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "tfb_query");
b.shader->info.cs.local_size[0] = 64;
b.shader->info.cs.local_size[1] = 1;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 64;
b.shader->info.cs.workgroup_size[1] = 1;
b.shader->info.cs.workgroup_size[2] = 1;
/* Create and initialize local variables. */
nir_variable *result =
@ -443,8 +443,8 @@ build_tfb_query_shader(struct radv_device *device)
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.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.
@ -552,9 +552,9 @@ build_timestamp_query_shader(struct radv_device *device)
* }
*/
nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "timestamp_query");
b.shader->info.cs.local_size[0] = 64;
b.shader->info.cs.local_size[1] = 1;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 64;
b.shader->info.cs.workgroup_size[1] = 1;
b.shader->info.cs.workgroup_size[2] = 1;
/* Create and initialize local variables. */
nir_variable *result = nir_local_variable_create(b.impl, glsl_uint64_t_type(), "result");
@ -573,8 +573,8 @@ build_timestamp_query_shader(struct radv_device *device)
nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *block_size =
nir_imm_ivec4(&b, b.shader->info.cs.local_size[0], b.shader->info.cs.local_size[1],
b.shader->info.cs.local_size[2], 0);
nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
b.shader->info.cs.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.

2
src/amd/vulkan/radv_shader_info.c
View File

@ -649,7 +649,7 @@ radv_nir_shader_info_pass(struct radv_device *device, const struct nir_shader *n
switch (nir->info.stage) {
case MESA_SHADER_COMPUTE:
for (int i = 0; i < 3; ++i)
info->cs.block_size[i] = nir->info.cs.local_size[i];
info->cs.block_size[i] = nir->info.cs.workgroup_size[i];
break;
case MESA_SHADER_FRAGMENT:
info->ps.can_discard = nir->info.fs.uses_discard;

12
src/broadcom/compiler/nir_to_vir.c
View File

@ -3744,9 +3744,9 @@ nir_to_vir(struct v3d_compile *c)
/* Set up the division between gl_LocalInvocationIndex and
* wg_in_mem in the payload reg.
*/
int wg_size = (c->s->info.cs.local_size[0] *
c->s->info.cs.local_size[1] *
c->s->info.cs.local_size[2]);
int wg_size = (c->s->info.cs.workgroup_size[0] *
c->s->info.cs.workgroup_size[1] *
c->s->info.cs.workgroup_size[2]);
c->local_invocation_index_bits =
ffs(util_next_power_of_two(MAX2(wg_size, 64))) - 1;
assert(c->local_invocation_index_bits <= 8);
@ -3754,9 +3754,9 @@ nir_to_vir(struct v3d_compile *c)
if (c->s->info.shared_size) {
struct qreg wg_in_mem = vir_SHR(c, c->cs_payload[1],
vir_uniform_ui(c, 16));
if (c->s->info.cs.local_size[0] != 1 ||
c->s->info.cs.local_size[1] != 1 ||
c->s->info.cs.local_size[2] != 1) {
if (c->s->info.cs.workgroup_size[0] != 1 ||
c->s->info.cs.workgroup_size[1] != 1 ||
c->s->info.cs.workgroup_size[2] != 1) {
int wg_bits = (16 -
c->local_invocation_index_bits);
int wg_mask = (1 << wg_bits) - 1;

6
src/broadcom/compiler/vir.c
View File

@ -804,9 +804,9 @@ v3d_cs_set_prog_data(struct v3d_compile *c,
{
prog_data->shared_size = c->s->info.shared_size;
prog_data->local_size[0] = c->s->info.cs.local_size[0];
prog_data->local_size[1] = c->s->info.cs.local_size[1];
prog_data->local_size[2] = c->s->info.cs.local_size[2];
prog_data->local_size[0] = c->s->info.cs.workgroup_size[0];
prog_data->local_size[1] = c->s->info.cs.workgroup_size[1];
prog_data->local_size[2] = c->s->info.cs.workgroup_size[2];
}
static void

28
src/compiler/glsl/linker.cpp
View File

@ -2230,9 +2230,9 @@ link_cs_input_layout_qualifiers(struct gl_shader_program *prog,
return;
for (int i = 0; i < 3; i++)
gl_prog->info.cs.local_size[i] = 0;
gl_prog->info.cs.workgroup_size[i] = 0;
gl_prog->info.cs.local_size_variable = false;
gl_prog->info.cs.workgroup_size_variable = false;
gl_prog->info.cs.derivative_group = DERIVATIVE_GROUP_NONE;
@ -2250,9 +2250,9 @@ link_cs_input_layout_qualifiers(struct gl_shader_program *prog,
struct gl_shader *shader = shader_list[sh];
if (shader->info.Comp.LocalSize[0] != 0) {
if (gl_prog->info.cs.local_size[0] != 0) {
if (gl_prog->info.cs.workgroup_size[0] != 0) {
for (int i = 0; i < 3; i++) {
if (gl_prog->info.cs.local_size[i] !=
if (gl_prog->info.cs.workgroup_size[i] !=
shader->info.Comp.LocalSize[i]) {
linker_error(prog, "compute shader defined with conflicting "
"local sizes\n");
@ -2261,11 +2261,11 @@ link_cs_input_layout_qualifiers(struct gl_shader_program *prog,
}
}
for (int i = 0; i < 3; i++) {
gl_prog->info.cs.local_size[i] =
gl_prog->info.cs.workgroup_size[i] =
shader->info.Comp.LocalSize[i];
}
} else if (shader->info.Comp.LocalSizeVariable) {
if (gl_prog->info.cs.local_size[0] != 0) {
if (gl_prog->info.cs.workgroup_size[0] != 0) {
/* The ARB_compute_variable_group_size spec says:
*
* If one compute shader attached to a program declares a
@ -2277,7 +2277,7 @@ link_cs_input_layout_qualifiers(struct gl_shader_program *prog,
"variable local group size\n");
return;
}
gl_prog->info.cs.local_size_variable = true;
gl_prog->info.cs.workgroup_size_variable = true;
}
enum gl_derivative_group group = shader->info.Comp.DerivativeGroup;
@ -2296,30 +2296,30 @@ link_cs_input_layout_qualifiers(struct gl_shader_program *prog,
* since we already know we're in the right type of shader program
* for doing it.
*/
if (gl_prog->info.cs.local_size[0] == 0 &&
!gl_prog->info.cs.local_size_variable) {
if (gl_prog->info.cs.workgroup_size[0] == 0 &&
!gl_prog->info.cs.workgroup_size_variable) {
linker_error(prog, "compute shader must contain a fixed or a variable "
"local group size\n");
return;
}
if (gl_prog->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS) {
if (gl_prog->info.cs.local_size[0] % 2 != 0) {
if (gl_prog->info.cs.workgroup_size[0] % 2 != 0) {
linker_error(prog, "derivative_group_quadsNV must be used with a "
"local group size whose first dimension "
"is a multiple of 2\n");
return;
}
if (gl_prog->info.cs.local_size[1] % 2 != 0) {
if (gl_prog->info.cs.workgroup_size[1] % 2 != 0) {
linker_error(prog, "derivative_group_quadsNV must be used with a local"
"group size whose second dimension "
"is a multiple of 2\n");
return;
}
} else if (gl_prog->info.cs.derivative_group == DERIVATIVE_GROUP_LINEAR) {
if ((gl_prog->info.cs.local_size[0] *
gl_prog->info.cs.local_size[1] *
gl_prog->info.cs.local_size[2]) % 4 != 0) {
if ((gl_prog->info.cs.workgroup_size[0] *
gl_prog->info.cs.workgroup_size[1] *
gl_prog->info.cs.workgroup_size[2]) % 4 != 0) {
linker_error(prog, "derivative_group_linearNV must be used with a "
"local group size whose total number of invocations "
"is a multiple of 4\n");

4
src/compiler/glsl/lower_cs_derived.cpp
View File

@ -54,7 +54,7 @@ public:
explicit lower_cs_derived_visitor(gl_linked_shader *shader)
: progress(false),
shader(shader),
local_size_variable(shader->Program->info.cs.local_size_variable),
local_size_variable(shader->Program->info.cs.workgroup_size_variable),
gl_WorkGroupSize(NULL),
gl_WorkGroupID(NULL),
gl_LocalInvocationID(NULL),
@ -144,7 +144,7 @@ lower_cs_derived_visitor::find_sysvals()
ir_constant_data data;
memset(&data, 0, sizeof(data));
for (int i = 0; i < 3; i++)
data.u[i] = shader->Program->info.cs.local_size[i];
data.u[i] = shader->Program->info.cs.workgroup_size[i];
gl_WorkGroupSize = new(shader) ir_constant(glsl_type::uvec3_type, &data);
}
}

24
src/compiler/nir/nir_lower_system_values.c
View File

@ -320,10 +320,10 @@ lower_compute_system_value_instr(nir_builder *b,
nir_ssa_def *x = nir_channel(b, ids, 0);
nir_ssa_def *y = nir_channel(b, ids, 1);
nir_ssa_def *z = nir_channel(b, ids, 2);
unsigned size_x = b->shader->info.cs.local_size[0];
unsigned size_x = b->shader->info.cs.workgroup_size[0];
nir_ssa_def *size_x_imm;
if (b->shader->info.cs.local_size_variable)
if (b->shader->info.cs.workgroup_size_variable)
size_x_imm = nir_channel(b, nir_load_local_group_size(b), 0);
else
size_x_imm = nir_imm_int(b, size_x);
@ -371,7 +371,7 @@ lower_compute_system_value_instr(nir_builder *b,
nir_ishl(b, x_bits_1n, one));
nir_ssa_def *i;
if (!b->shader->info.cs.local_size_variable &&
if (!b->shader->info.cs.workgroup_size_variable &&
util_is_power_of_two_nonzero(size_x)) {
nir_ssa_def *log2_size_x = nir_imm_int(b, util_logbase2(size_x));
i = nir_ior(b, bits_01x, nir_ishl(b, y_bits_1n, log2_size_x));
@ -405,9 +405,9 @@ lower_compute_system_value_instr(nir_builder *b,
nir_ssa_def *local_id = nir_load_local_invocation_id(b);
nir_ssa_def *size_x =
nir_imm_int(b, b->shader->info.cs.local_size[0]);
nir_imm_int(b, b->shader->info.cs.workgroup_size[0]);
nir_ssa_def *size_y =
nir_imm_int(b, b->shader->info.cs.local_size[1]);
nir_imm_int(b, b->shader->info.cs.workgroup_size[1]);
/* Because no hardware supports a local workgroup size greater than
* about 1K, this calculation can be done in 32-bit and can save some
@ -425,7 +425,7 @@ lower_compute_system_value_instr(nir_builder *b,
}
case nir_intrinsic_load_local_group_size:
if (b->shader->info.cs.local_size_variable) {
if (b->shader->info.cs.workgroup_size_variable) {
/* If the local work group size is variable it can't be lowered at
* this point. We do, however, have to make sure that the intrinsic
* is only 32-bit.
@ -434,12 +434,12 @@ lower_compute_system_value_instr(nir_builder *b,
} else {
/* using a 32 bit constant is safe here as no device/driver needs more
* than 32 bits for the local size */
nir_const_value local_size_const[3];
memset(local_size_const, 0, sizeof(local_size_const));
local_size_const[0].u32 = b->shader->info.cs.local_size[0];
local_size_const[1].u32 = b->shader->info.cs.local_size[1];
local_size_const[2].u32 = b->shader->info.cs.local_size[2];
return nir_u2u(b, nir_build_imm(b, 3, 32, local_size_const), bit_size);
nir_const_value workgroup_size_const[3];
memset(workgroup_size_const, 0, sizeof(workgroup_size_const));
workgroup_size_const[0].u32 = b->shader->info.cs.workgroup_size[0];
workgroup_size_const[1].u32 = b->shader->info.cs.workgroup_size[1];
workgroup_size_const[2].u32 = b->shader->info.cs.workgroup_size[2];
return nir_u2u(b, nir_build_imm(b, 3, 32, workgroup_size_const), bit_size);
}
case nir_intrinsic_load_global_invocation_id_zero_base: {

8
src/compiler/nir/nir_lower_variable_initializers.c
View File

@ -154,10 +154,10 @@ nir_zero_initialize_shared_memory(nir_shader *shader,
nir_builder_init(&b, nir_shader_get_entrypoint(shader));
b.cursor = nir_before_cf_list(&b.impl->body);
assert(!shader->info.cs.local_size_variable);
const unsigned local_count = shader->info.cs.local_size[0] *
shader->info.cs.local_size[1] *
shader->info.cs.local_size[2];
assert(!shader->info.cs.workgroup_size_variable);
const unsigned local_count = shader->info.cs.workgroup_size[0] *
shader->info.cs.workgroup_size[1] *
shader->info.cs.workgroup_size[2];
/* The initialization logic is simplified if we can always split the memory
* in full chunk_size units.

8
src/compiler/nir/nir_opt_uniform_atomics.c
View File

@ -169,7 +169,7 @@ is_atomic_already_optimized(nir_shader *shader, nir_intrinsic_instr *instr)
unsigned dims_needed = 0;
for (unsigned i = 0; i < 3; i++)
dims_needed |= (shader->info.cs.local_size[i] > 1) << i;
dims_needed |= (shader->info.cs.workgroup_size[i] > 1) << i;
return (dims & dims_needed) == dims_needed || dims & 0x8;
}
@ -306,9 +306,9 @@ nir_opt_uniform_atomics(nir_shader *shader)
/* A 1x1x1 workgroup only ever has one active lane, so there's no point in
* optimizing any atomics.
*/
if (shader->info.stage == MESA_SHADER_COMPUTE && !shader->info.cs.local_size_variable &&
shader->info.cs.local_size[0] == 1 && shader->info.cs.local_size[1] == 1 &&
shader->info.cs.local_size[2] == 1)
if (shader->info.stage == MESA_SHADER_COMPUTE && !shader->info.cs.workgroup_size_variable &&
shader->info.cs.workgroup_size[0] == 1 && shader->info.cs.workgroup_size[1] == 1 &&
shader->info.cs.workgroup_size[2] == 1)
return false;
nir_foreach_function(function, shader) {

10
src/compiler/nir/nir_print.c
View File

@ -1606,11 +1606,11 @@ nir_print_shader_annotated(nir_shader *shader, FILE *fp,
fprintf(fp, "label: %s\n", shader->info.label);
if (gl_shader_stage_is_compute(shader->info.stage)) {
fprintf(fp, "local-size: %u, %u, %u%s\n",
shader->info.cs.local_size[0],
shader->info.cs.local_size[1],
shader->info.cs.local_size[2],
shader->info.cs.local_size_variable ? " (variable)" : "");
fprintf(fp, "workgroup-size: %u, %u, %u%s\n",
shader->info.cs.workgroup_size[0],
shader->info.cs.workgroup_size[1],
shader->info.cs.workgroup_size[2],
shader->info.cs.workgroup_size_variable ? " (variable)" : "");
fprintf(fp, "shared-size: %u\n", shader->info.shared_size);
}

24
src/compiler/nir/nir_range_analysis.c
View File

@ -1293,19 +1293,19 @@ nir_unsigned_upper_bound(nir_shader *shader, struct hash_table *range_ht,
switch (intrin->intrinsic) {
case nir_intrinsic_load_local_invocation_index:
if (shader->info.stage != MESA_SHADER_COMPUTE ||
shader->info.cs.local_size_variable) {
shader->info.cs.workgroup_size_variable) {
res = config->max_work_group_invocations - 1;
} else {
res = (shader->info.cs.local_size[0] *
shader->info.cs.local_size[1] *
shader->info.cs.local_size[2]) - 1u;
res = (shader->info.cs.workgroup_size[0] *
shader->info.cs.workgroup_size[1] *
shader->info.cs.workgroup_size[2]) - 1u;
}
break;
case nir_intrinsic_load_local_invocation_id:
if (shader->info.cs.local_size_variable)
if (shader->info.cs.workgroup_size_variable)
res = config->max_work_group_size[scalar.comp] - 1u;
else
res = shader->info.cs.local_size[scalar.comp] - 1u;
res = shader->info.cs.workgroup_size[scalar.comp] - 1u;
break;
case nir_intrinsic_load_work_group_id:
res = config->max_work_group_count[scalar.comp] - 1u;
@ -1314,11 +1314,11 @@ nir_unsigned_upper_bound(nir_shader *shader, struct hash_table *range_ht,
res = config->max_work_group_count[scalar.comp];
break;
case nir_intrinsic_load_global_invocation_id:
if (shader->info.cs.local_size_variable) {
if (shader->info.cs.workgroup_size_variable) {
res = mul_clamp(config->max_work_group_size[scalar.comp],
config->max_work_group_count[scalar.comp]) - 1u;
} else {
res = (shader->info.cs.local_size[scalar.comp] *
res = (shader->info.cs.workgroup_size[scalar.comp] *
config->max_work_group_count[scalar.comp]) - 1u;
}
break;
@ -1339,10 +1339,10 @@ nir_unsigned_upper_bound(nir_shader *shader, struct hash_table *range_ht,
case nir_intrinsic_load_subgroup_id:
case nir_intrinsic_load_num_subgroups: {
uint32_t work_group_size = config->max_work_group_invocations;
if (shader->info.stage == MESA_SHADER_COMPUTE && !shader->info.cs.local_size_variable) {
work_group_size = shader->info.cs.local_size[0] *
shader->info.cs.local_size[1] *
shader->info.cs.local_size[2];
if (shader->info.stage == MESA_SHADER_COMPUTE && !shader->info.cs.workgroup_size_variable) {
work_group_size = shader->info.cs.workgroup_size[0] *
shader->info.cs.workgroup_size[1] *
shader->info.cs.workgroup_size[2];
}
res = DIV_ROUND_UP(work_group_size, config->min_subgroup_size);
if (intrin->intrinsic == nir_intrinsic_load_subgroup_id)

6
src/compiler/shader_info.h
View File

@ -384,10 +384,10 @@ typedef struct shader_info {
} fs;
struct {
uint16_t local_size[3];
uint16_t local_size_hint[3];
uint16_t workgroup_size[3];
uint16_t workgroup_size_hint[3];
bool local_size_variable:1;
bool workgroup_size_variable:1;
uint8_t user_data_components_amd:3;
/*

30
src/compiler/spirv/spirv_to_nir.c
View File

@ -4796,16 +4796,16 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
case SpvExecutionModeLocalSizeHint:
vtn_assert(b->shader->info.stage == MESA_SHADER_KERNEL);
b->shader->info.cs.local_size_hint[0] = mode->operands[0];
b->shader->info.cs.local_size_hint[1] = mode->operands[1];
b->shader->info.cs.local_size_hint[2] = mode->operands[2];
b->shader->info.cs.workgroup_size_hint[0] = mode->operands[0];
b->shader->info.cs.workgroup_size_hint[1] = mode->operands[1];
b->shader->info.cs.workgroup_size_hint[2] = mode->operands[2];
break;
case SpvExecutionModeLocalSize:
vtn_assert(gl_shader_stage_is_compute(b->shader->info.stage));
b->shader->info.cs.local_size[0] = mode->operands[0];
b->shader->info.cs.local_size[1] = mode->operands[1];
b->shader->info.cs.local_size[2] = mode->operands[2];
b->shader->info.cs.workgroup_size[0] = mode->operands[0];
b->shader->info.cs.workgroup_size[1] = mode->operands[1];
b->shader->info.cs.workgroup_size[2] = mode->operands[2];
break;
case SpvExecutionModeOutputVertices:
@ -5016,16 +5016,16 @@ vtn_handle_execution_mode_id(struct vtn_builder *b, struct vtn_value *entry_poin
switch (mode->exec_mode) {
case SpvExecutionModeLocalSizeId:
b->shader->info.cs.local_size[0] = vtn_constant_uint(b, mode->operands[0]);
b->shader->info.cs.local_size[1] = vtn_constant_uint(b, mode->operands[1]);
b->shader->info.cs.local_size[2] = vtn_constant_uint(b, mode->operands[2]);
b->shader->info.cs.workgroup_size[0] = vtn_constant_uint(b, mode->operands[0]);
b->shader->info.cs.workgroup_size[1] = vtn_constant_uint(b, mode->operands[1]);
b->shader->info.cs.workgroup_size[2] = vtn_constant_uint(b, mode->operands[2]);
break;
case SpvExecutionModeLocalSizeHintId:
vtn_assert(b->shader->info.stage == MESA_SHADER_KERNEL);
b->shader->info.cs.local_size_hint[0] = vtn_constant_uint(b, mode->operands[0]);
b->shader->info.cs.local_size_hint[1] = vtn_constant_uint(b, mode->operands[1]);
b->shader->info.cs.local_size_hint[2] = vtn_constant_uint(b, mode->operands[2]);
b->shader->info.cs.workgroup_size_hint[0] = vtn_constant_uint(b, mode->operands[0]);
b->shader->info.cs.workgroup_size_hint[1] = vtn_constant_uint(b, mode->operands[1]);
b->shader->info.cs.workgroup_size_hint[2] = vtn_constant_uint(b, mode->operands[2]);
break;
default:
@ -5993,9 +5993,9 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
nir_const_value *const_size =
b->workgroup_size_builtin->constant->values;
b->shader->info.cs.local_size[0] = const_size[0].u32;
b->shader->info.cs.local_size[1] = const_size[1].u32;
b->shader->info.cs.local_size[2] = const_size[2].u32;
b->shader->info.cs.workgroup_size[0] = const_size[0].u32;
b->shader->info.cs.workgroup_size[1] = const_size[1].u32;
b->shader->info.cs.workgroup_size[2] = const_size[2].u32;
}
/* Set types on all vtn_values */

8
src/freedreno/ir3/ir3_compiler_nir.c
View File

@ -4102,10 +4102,10 @@ ir3_compile_shader_nir(struct ir3_compiler *compiler,
so->need_pixlod = true;
if (so->type == MESA_SHADER_COMPUTE) {
so->local_size[0] = ctx->s->info.cs.local_size[0];
so->local_size[1] = ctx->s->info.cs.local_size[1];
so->local_size[2] = ctx->s->info.cs.local_size[2];
so->local_size_variable = ctx->s->info.cs.local_size_variable;
so->local_size[0] = ctx->s->info.cs.workgroup_size[0];
so->local_size[1] = ctx->s->info.cs.workgroup_size[1];
so->local_size[2] = ctx->s->info.cs.workgroup_size[2];
so->local_size_variable = ctx->s->info.cs.workgroup_size_variable;
}
out:

6
src/gallium/auxiliary/nir/nir_to_tgsi_info.c
View File

@ -487,9 +487,9 @@ void nir_tgsi_scan_shader(const struct nir_shader *nir,
}
if (gl_shader_stage_is_compute(nir->info.stage)) {
info->properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] = nir->info.cs.local_size[0];
info->properties[TGSI_PROPERTY_CS_FIXED_BLOCK_HEIGHT] = nir->info.cs.local_size[1];
info->properties[TGSI_PROPERTY_CS_FIXED_BLOCK_DEPTH] = nir->info.cs.local_size[2];
info->properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] = nir->info.cs.workgroup_size[0];
info->properties[TGSI_PROPERTY_CS_FIXED_BLOCK_HEIGHT] = nir->info.cs.workgroup_size[1];
info->properties[TGSI_PROPERTY_CS_FIXED_BLOCK_DEPTH] = nir->info.cs.workgroup_size[2];
}
i = 0;

14
src/gallium/auxiliary/nir/tgsi_to_nir.c
View File

@ -2354,15 +2354,15 @@ ttn_compile_init(const void *tgsi_tokens,
break;
case TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH:
if (s->info.stage == MESA_SHADER_COMPUTE)
s->info.cs.local_size[0] = value;
s->info.cs.workgroup_size[0] = value;
break;
case TGSI_PROPERTY_CS_FIXED_BLOCK_HEIGHT:
if (s->info.stage == MESA_SHADER_COMPUTE)
s->info.cs.local_size[1] = value;
s->info.cs.workgroup_size[1] = value;
break;
case TGSI_PROPERTY_CS_FIXED_BLOCK_DEPTH:
if (s->info.stage == MESA_SHADER_COMPUTE)
s->info.cs.local_size[2] = value;
s->info.cs.workgroup_size[2] = value;
break;
case TGSI_PROPERTY_CS_USER_DATA_COMPONENTS_AMD:
if (s->info.stage == MESA_SHADER_COMPUTE)
@ -2381,10 +2381,10 @@ ttn_compile_init(const void *tgsi_tokens,
}
if (s->info.stage == MESA_SHADER_COMPUTE &&
(!s->info.cs.local_size[0] ||
!s->info.cs.local_size[1] ||
!s->info.cs.local_size[2]))
s->info.cs.local_size_variable = true;
(!s->info.cs.workgroup_size[0] ||
!s->info.cs.workgroup_size[1] ||
!s->info.cs.workgroup_size[2]))
s->info.cs.workgroup_size_variable = true;
c->inputs = rzalloc_array(c, struct nir_variable *, s->num_inputs);
c->outputs = rzalloc_array(c, struct nir_variable *, s->num_outputs);

6
src/gallium/auxiliary/tgsi/tgsi_ureg.c
View File

@ -2354,11 +2354,11 @@ ureg_setup_compute_shader(struct ureg_program *ureg,
const struct shader_info *info)
{
ureg_property(ureg, TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH,
info->cs.local_size[0]);
info->cs.workgroup_size[0]);
ureg_property(ureg, TGSI_PROPERTY_CS_FIXED_BLOCK_HEIGHT,
info->cs.local_size[1]);
info->cs.workgroup_size[1]);
ureg_property(ureg, TGSI_PROPERTY_CS_FIXED_BLOCK_DEPTH,
info->cs.local_size[2]);
info->cs.workgroup_size[2]);
if (info->shared_size)
ureg_DECL_memory(ureg, TGSI_MEMORY_TYPE_SHARED);

2
src/gallium/drivers/freedreno/a5xx/fd5_compute.c
View File

@ -142,7 +142,7 @@ fd5_launch_grid(struct fd_context *ctx,
}
const unsigned *local_size =
info->block; // v->shader->nir->info->cs.local_size;
info->block; // v->shader->nir->info->cs.workgroup_size;
const unsigned *num_groups = info->grid;
/* for some reason, mesa/st doesn't set info->work_dim, so just assume 3: */
const unsigned work_dim = info->work_dim ? info->work_dim : 3;

2
src/gallium/drivers/freedreno/a6xx/fd6_compute.c
View File

@ -134,7 +134,7 @@ fd6_launch_grid(struct fd_context *ctx, const struct pipe_grid_info *info) in_dt
OUT_RING(ring, A6XX_CP_SET_MARKER_0_MODE(RM6_COMPUTE));
const unsigned *local_size =
info->block; // v->shader->nir->info->cs.local_size;
info->block; // v->shader->nir->info->cs.workgroup_size;
const unsigned *num_groups = info->grid;
/* for some reason, mesa/st doesn't set info->work_dim, so just assume 3: */
const unsigned work_dim = info->work_dim ? info->work_dim : 3;

2
src/gallium/drivers/iris/iris_program.c
View File

@ -528,7 +528,7 @@ iris_setup_uniforms(const struct brw_compiler *compiler,
break;
}
case nir_intrinsic_load_local_group_size: {
assert(nir->info.cs.local_size_variable);
assert(nir->info.cs.workgroup_size_variable);
if (variable_group_size_idx == -1) {
variable_group_size_idx = num_system_values;
num_system_values += 3;

6
src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
View File

@ -1289,9 +1289,9 @@ Converter::parseNIR()
switch(prog->getType()) {
case Program::TYPE_COMPUTE:
info->prop.cp.numThreads[0] = nir->info.cs.local_size[0];
info->prop.cp.numThreads[1] = nir->info.cs.local_size[1];
info->prop.cp.numThreads[2] = nir->info.cs.local_size[2];
info->prop.cp.numThreads[0] = nir->info.cs.workgroup_size[0];
info->prop.cp.numThreads[1] = nir->info.cs.workgroup_size[1];
info->prop.cp.numThreads[2] = nir->info.cs.workgroup_size[2];
info_out->bin.smemSize += nir->info.shared_size;
break;
case Program::TYPE_FRAGMENT:

4
src/gallium/drivers/radeonsi/si_shader.c
View File

@ -218,10 +218,10 @@ unsigned si_get_max_workgroup_size(const struct si_shader *shader)
}
/* Compile a variable block size using the maximum variable size. */
if (shader->selector->info.base.cs.local_size_variable)
if (shader->selector->info.base.cs.workgroup_size_variable)
return SI_MAX_VARIABLE_THREADS_PER_BLOCK;
uint16_t *local_size = shader->selector->info.base.cs.local_size;
uint16_t *local_size = shader->selector->info.base.cs.workgroup_size;
unsigned max_work_group_size = (uint32_t)local_size[0] *
(uint32_t)local_size[1] *
(uint32_t)local_size[2];

2
src/gallium/drivers/radeonsi/si_shader_llvm.c
View File

@ -411,7 +411,7 @@ static LLVMValueRef si_llvm_get_block_size(struct ac_shader_abi *abi)
{
struct si_shader_context *ctx = si_shader_context_from_abi(abi);
assert(ctx->shader->selector->info.base.cs.local_size_variable &&
assert(ctx->shader->selector->info.base.cs.workgroup_size_variable &&
ctx->shader->selector->info.uses_variable_block_size);
LLVMValueRef chan[3] = {

12
src/gallium/drivers/radeonsi/si_shaderlib_nir.c
View File

@ -59,9 +59,9 @@ void *si_create_dcc_retile_cs(struct si_context *sctx, struct radeon_surf *surf)
sctx->b.screen->get_compiler_options(sctx->b.screen, PIPE_SHADER_IR_NIR, PIPE_SHADER_COMPUTE);
nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, options, "dcc_retile");
b.shader->info.cs.local_size[0] = 8;
b.shader->info.cs.local_size[1] = 8;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 8;
b.shader->info.cs.workgroup_size[1] = 8;
b.shader->info.cs.workgroup_size[2] = 1;
b.shader->info.cs.user_data_components_amd = 3;
b.shader->info.num_ssbos = 1;
@ -107,9 +107,9 @@ void *gfx9_create_clear_dcc_msaa_cs(struct si_context *sctx, struct si_texture *
sctx->b.screen->get_compiler_options(sctx->b.screen, PIPE_SHADER_IR_NIR, PIPE_SHADER_COMPUTE);
nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, options, "clear_dcc_msaa");
b.shader->info.cs.local_size[0] = 8;
b.shader->info.cs.local_size[1] = 8;
b.shader->info.cs.local_size[2] = 1;
b.shader->info.cs.workgroup_size[0] = 8;
b.shader->info.cs.workgroup_size[1] = 8;
b.shader->info.cs.workgroup_size[2] = 1;
b.shader->info.cs.user_data_components_amd = 2;
b.shader->info.num_ssbos = 1;

6
src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c
View File

@ -3806,10 +3806,10 @@ nir_to_spirv(struct nir_shader *s, const struct zink_so_info *so_info, uint32_t
if (s->info.shared_size)
create_shared_block(&ctx, s->info.shared_size);
if (s->info.cs.local_size[0] || s->info.cs.local_size[1] || s->info.cs.local_size[2])
if (s->info.cs.workgroup_size[0] || s->info.cs.workgroup_size[1] || s->info.cs.workgroup_size[2])
spirv_builder_emit_exec_mode_literal3(&ctx.builder, entry_point, SpvExecutionModeLocalSize,
(uint32_t[3]){(uint32_t)s->info.cs.local_size[0], (uint32_t)s->info.cs.local_size[1],
(uint32_t)s->info.cs.local_size[2]});
(uint32_t[3]){(uint32_t)s->info.cs.workgroup_size[0], (uint32_t)s->info.cs.workgroup_size[1],
(uint32_t)s->info.cs.workgroup_size[2]});
else {
SpvId sizes[3];
uint32_t ids[] = {ZINK_WORKGROUP_SIZE_X, ZINK_WORKGROUP_SIZE_Y, ZINK_WORKGROUP_SIZE_Z};

6
src/gallium/drivers/zink/zink_program.c
View File

@ -589,9 +589,9 @@ void
zink_program_update_compute_pipeline_state(struct zink_context *ctx, struct zink_compute_program *comp, const uint block[3])
{
struct zink_shader *zs = comp->shader;
bool use_local_size = !(zs->nir->info.cs.local_size[0] ||
zs->nir->info.cs.local_size[1] ||
zs->nir->info.cs.local_size[2]);
bool use_local_size = !(zs->nir->info.cs.workgroup_size[0] ||
zs->nir->info.cs.workgroup_size[1] ||
zs->nir->info.cs.workgroup_size[2]);
if (ctx->compute_pipeline_state.use_local_size != use_local_size)
ctx->compute_pipeline_state.dirty = true;
ctx->compute_pipeline_state.use_local_size = use_local_size;

8
src/gallium/frontends/clover/nir/invocation.cpp
View File

@ -428,10 +428,10 @@ module clover::nir::spirv_to_nir(const module &mod, const device &dev,
throw build_error();
}
nir->info.cs.local_size_variable = sym.reqd_work_group_size[0] == 0;
nir->info.cs.local_size[0] = sym.reqd_work_group_size[0];
nir->info.cs.local_size[1] = sym.reqd_work_group_size[1];
nir->info.cs.local_size[2] = sym.reqd_work_group_size[2];
nir->info.cs.workgroup_size_variable = sym.reqd_work_group_size[0] == 0;
nir->info.cs.workgroup_size[0] = sym.reqd_work_group_size[0];
nir->info.cs.workgroup_size[1] = sym.reqd_work_group_size[1];
nir->info.cs.workgroup_size[2] = sym.reqd_work_group_size[2];
nir_validate_shader(nir, "clover");
// Inline all functions first.

6
src/gallium/frontends/lavapipe/lvp_execute.c
View File

@ -348,9 +348,9 @@ static void handle_compute_pipeline(struct lvp_cmd_buffer_entry *cmd,
{
struct lvp_pipeline *pipeline = cmd->u.pipeline.pipeline;
state->dispatch_info.block[0] = pipeline->pipeline_nir[MESA_SHADER_COMPUTE]->info.cs.local_size[0];
state->dispatch_info.block[1] = pipeline->pipeline_nir[MESA_SHADER_COMPUTE]->info.cs.local_size[1];
state->dispatch_info.block[2] = pipeline->pipeline_nir[MESA_SHADER_COMPUTE]->info.cs.local_size[2];
state->dispatch_info.block[0] = pipeline->pipeline_nir[MESA_SHADER_COMPUTE]->info.cs.workgroup_size[0];
state->dispatch_info.block[1] = pipeline->pipeline_nir[MESA_SHADER_COMPUTE]->info.cs.workgroup_size[1];
state->dispatch_info.block[2] = pipeline->pipeline_nir[MESA_SHADER_COMPUTE]->info.cs.workgroup_size[2];
state->pctx->bind_compute_state(state->pctx, pipeline->shader_cso[PIPE_SHADER_COMPUTE]);
}

16
src/intel/compiler/brw_fs.cpp
View File

@ -9534,10 +9534,10 @@ lower_simd(nir_builder *b, nir_instr *instr, void *options)
/* If the whole workgroup fits in one thread, we can lower subgroup_id
* to a constant zero.
*/
if (!b->shader->info.cs.local_size_variable) {
unsigned local_workgroup_size = b->shader->info.cs.local_size[0] *
b->shader->info.cs.local_size[1] *
b->shader->info.cs.local_size[2];
if (!b->shader->info.cs.workgroup_size_variable) {
unsigned local_workgroup_size = b->shader->info.cs.workgroup_size[0] *
b->shader->info.cs.workgroup_size[1] *
b->shader->info.cs.workgroup_size[2];
if (local_workgroup_size <= simd_width)
return nir_imm_int(b, 0);
}
@ -9599,15 +9599,15 @@ brw_compile_cs(const struct brw_compiler *compiler,
unsigned min_dispatch_width;
unsigned max_dispatch_width;
if (nir->info.cs.local_size_variable) {
if (nir->info.cs.workgroup_size_variable) {
generate_all = true;
min_dispatch_width = 8;
max_dispatch_width = 32;
} else {
generate_all = false;
prog_data->local_size[0] = nir->info.cs.local_size[0];
prog_data->local_size[1] = nir->info.cs.local_size[1];
prog_data->local_size[2] = nir->info.cs.local_size[2];
prog_data->local_size[0] = nir->info.cs.workgroup_size[0];
prog_data->local_size[1] = nir->info.cs.workgroup_size[1];
prog_data->local_size[2] = nir->info.cs.workgroup_size[2];
unsigned local_workgroup_size = prog_data->local_size[0] *
prog_data->local_size[1] *
prog_data->local_size[2];

8
src/intel/compiler/brw_fs_nir.cpp
View File

@ -113,7 +113,7 @@ fs_visitor::nir_setup_uniforms()
assert(uniforms == prog_data->nr_params);
uint32_t *param;
if (nir->info.cs.local_size_variable &&
if (nir->info.cs.workgroup_size_variable &&
compiler->lower_variable_group_size) {
param = brw_stage_prog_data_add_params(prog_data, 3);
for (unsigned i = 0; i < 3; i++) {
@ -3671,7 +3671,7 @@ fs_visitor::nir_emit_cs_intrinsic(const fs_builder &bld,
* invocations are already executed lock-step. Instead of an actual
* barrier just emit a scheduling fence, that will generate no code.
*/
if (!nir->info.cs.local_size_variable &&
if (!nir->info.cs.workgroup_size_variable &&
workgroup_size() <= dispatch_width) {
bld.exec_all().group(1, 0).emit(FS_OPCODE_SCHEDULING_FENCE);
break;
@ -3816,7 +3816,7 @@ fs_visitor::nir_emit_cs_intrinsic(const fs_builder &bld,
case nir_intrinsic_load_local_group_size: {
assert(compiler->lower_variable_group_size);
assert(nir->info.cs.local_size_variable);
assert(nir->info.cs.workgroup_size_variable);
for (unsigned i = 0; i < 3; i++) {
bld.MOV(retype(offset(dest, bld, i), BRW_REGISTER_TYPE_UD),
group_size[i]);
@ -4324,7 +4324,7 @@ fs_visitor::nir_emit_intrinsic(const fs_builder &bld, nir_intrinsic_instr *instr
*
* TODO: Check if applies for many HW threads sharing same Data Port.
*/
if (!nir->info.cs.local_size_variable &&
if (!nir->info.cs.workgroup_size_variable &&
slm_fence && workgroup_size() <= dispatch_width)
slm_fence = false;

34
src/intel/compiler/brw_nir_lower_cs_intrinsics.c
View File

@ -81,13 +81,13 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state,
nir_ssa_def *size_x;
nir_ssa_def *size_y;
if (state->nir->info.cs.local_size_variable) {
if (state->nir->info.cs.workgroup_size_variable) {
nir_ssa_def *size_xyz = nir_load_local_group_size(b);
size_x = nir_channel(b, size_xyz, 0);
size_y = nir_channel(b, size_xyz, 1);
} else {
size_x = nir_imm_int(b, nir->info.cs.local_size[0]);
size_y = nir_imm_int(b, nir->info.cs.local_size[1]);
size_x = nir_imm_int(b, nir->info.cs.workgroup_size[0]);
size_y = nir_imm_int(b, nir->info.cs.workgroup_size[1]);
}
nir_ssa_def *size_xy = nir_imul(b, size_x, size_y);
@ -120,8 +120,8 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state,
id_x = nir_umod(b, linear, size_x);
id_y = nir_umod(b, nir_udiv(b, linear, size_x), size_y);
local_index = linear;
} else if (!nir->info.cs.local_size_variable &&
nir->info.cs.local_size[1] % 4 == 0) {
} else if (!nir->info.cs.workgroup_size_variable &&
nir->info.cs.workgroup_size[1] % 4 == 0) {
/* 1x4 block X-major lid order. Same as X-major except increments in
* blocks of width=1 height=4. Always optimal for tileY and usually
* optimal for linear accesses.
@ -213,16 +213,16 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state,
case nir_intrinsic_load_num_subgroups: {
nir_ssa_def *size;
if (state->nir->info.cs.local_size_variable) {
if (state->nir->info.cs.workgroup_size_variable) {
nir_ssa_def *size_xyz = nir_load_local_group_size(b);
nir_ssa_def *size_x = nir_channel(b, size_xyz, 0);
nir_ssa_def *size_y = nir_channel(b, size_xyz, 1);
nir_ssa_def *size_z = nir_channel(b, size_xyz, 2);
size = nir_imul(b, nir_imul(b, size_x, size_y), size_z);
} else {
size = nir_imm_int(b, nir->info.cs.local_size[0] *
nir->info.cs.local_size[1] *
nir->info.cs.local_size[2]);
size = nir_imm_int(b, nir->info.cs.workgroup_size[0] *
nir->info.cs.workgroup_size[1] *
nir->info.cs.workgroup_size[2]);
}
/* Calculate the equivalent of DIV_ROUND_UP. */
@ -273,16 +273,16 @@ brw_nir_lower_cs_intrinsics(nir_shader *nir)
};
/* Constraints from NV_compute_shader_derivatives. */
if (!nir->info.cs.local_size_variable) {
if (!nir->info.cs.workgroup_size_variable) {
if (nir->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS) {
assert(nir->info.cs.local_size[0] % 2 == 0);
assert(nir->info.cs.local_size[1] % 2 == 0);
assert(nir->info.cs.workgroup_size[0] % 2 == 0);
assert(nir->info.cs.workgroup_size[1] % 2 == 0);
} else if (nir->info.cs.derivative_group == DERIVATIVE_GROUP_LINEAR) {
ASSERTED unsigned local_workgroup_size =
nir->info.cs.local_size[0] *
nir->info.cs.local_size[1] *
nir->info.cs.local_size[2];
assert(local_workgroup_size % 4 == 0);
ASSERTED unsigned workgroup_size =
nir->info.cs.workgroup_size[0] *
nir->info.cs.workgroup_size[1] *
nir->info.cs.workgroup_size[2];
assert(workgroup_size % 4 == 0);
}
}

2
src/intel/compiler/brw_nir_rt.c
View File

@ -426,7 +426,7 @@ brw_nir_create_raygen_trampoline(const struct brw_compiler *compiler,
"RT Ray-Gen Trampoline");
ralloc_steal(mem_ctx, b.shader);
b.shader->info.cs.local_size_variable = true;
b.shader->info.cs.workgroup_size_variable = true;
/* The RT global data and raygen BINDLESS_SHADER_RECORD addresses are
* passed in as push constants in the first register. We deal with the

6
src/mesa/main/compute.c
View File

@ -89,7 +89,7 @@ validate_DispatchCompute(struct gl_context *ctx, const GLuint *num_groups)
* program for the compute shader stage has a variable work group size."
*/
struct gl_program *prog = ctx->_Shader->CurrentProgram[MESA_SHADER_COMPUTE];
if (prog->info.cs.local_size_variable) {
if (prog->info.cs.workgroup_size_variable) {
_mesa_error(ctx, GL_INVALID_OPERATION,
"glDispatchCompute(variable work group size forbidden)");
return GL_FALSE;
@ -113,7 +113,7 @@ validate_DispatchComputeGroupSizeARB(struct gl_context *ctx,
* shader stage has a fixed work group size."
*/
struct gl_program *prog = ctx->_Shader->CurrentProgram[MESA_SHADER_COMPUTE];
if (!prog->info.cs.local_size_variable) {
if (!prog->info.cs.workgroup_size_variable) {
_mesa_error(ctx, GL_INVALID_OPERATION,
"glDispatchComputeGroupSizeARB(fixed work group size "
"forbidden)");
@ -269,7 +269,7 @@ valid_dispatch_indirect(struct gl_context *ctx, GLintptr indirect)
* compute shader stage has a variable work group size."
*/
struct gl_program *prog = ctx->_Shader->CurrentProgram[MESA_SHADER_COMPUTE];
if (prog->info.cs.local_size_variable) {
if (prog->info.cs.workgroup_size_variable) {
_mesa_error(ctx, GL_INVALID_OPERATION,
"%s(variable work group size forbidden)", name);
return GL_FALSE;

2
src/mesa/main/shaderapi.c
View File

@ -935,7 +935,7 @@ get_programiv(struct gl_context *ctx, GLuint program, GLenum pname,
}
for (i = 0; i < 3; i++)
params[i] = shProg->_LinkedShaders[MESA_SHADER_COMPUTE]->
Program->info.cs.local_size[i];
Program->info.cs.workgroup_size[i];
return;
}
case GL_PROGRAM_SEPARABLE:

2
src/mesa/state_tracker/st_cb_compute.c
View File

@ -59,7 +59,7 @@ static void st_dispatch_compute_common(struct gl_context *ctx,
st_validate_state(st, ST_PIPELINE_COMPUTE);
for (unsigned i = 0; i < 3; i++) {
info.block[i] = group_size ? group_size[i] : prog->info.cs.local_size[i];
info.block[i] = group_size ? group_size[i] : prog->info.cs.workgroup_size[i];
info.grid[i] = num_groups ? num_groups[i] : 0;
}

24
src/microsoft/clc/clc_compiler.c
View File

@ -1080,7 +1080,7 @@ clc_to_dxil(struct clc_context *ctx,
clc_error(logger, "spirv_to_nir() failed");
goto err_free_dxil;
}
nir->info.cs.local_size_variable = true;
nir->info.cs.workgroup_size_variable = true;
NIR_PASS_V(nir, nir_lower_goto_ifs);
NIR_PASS_V(nir, nir_opt_dead_cf);
@ -1338,33 +1338,33 @@ clc_to_dxil(struct clc_context *ctx,
nir_variable *work_properties_var =
add_work_properties_var(dxil, nir, &cbv_id);
memcpy(metadata->local_size, nir->info.cs.local_size,
memcpy(metadata->local_size, nir->info.cs.workgroup_size,
sizeof(metadata->local_size));
memcpy(metadata->local_size_hint, nir->info.cs.local_size_hint,
memcpy(metadata->local_size_hint, nir->info.cs.workgroup_size_hint,
sizeof(metadata->local_size));
// Patch the localsize before calling clc_nir_lower_system_values().
if (conf) {
for (unsigned i = 0; i < ARRAY_SIZE(nir->info.cs.local_size); i++) {
for (unsigned i = 0; i < ARRAY_SIZE(nir->info.cs.workgroup_size); i++) {
if (!conf->local_size[i] ||
conf->local_size[i] == nir->info.cs.local_size[i])
conf->local_size[i] == nir->info.cs.workgroup_size[i])
continue;
if (nir->info.cs.local_size[i] &&
nir->info.cs.local_size[i] != conf->local_size[i]) {
if (nir->info.cs.workgroup_size[i] &&
nir->info.cs.workgroup_size[i] != conf->local_size[i]) {
debug_printf("D3D12: runtime local size does not match reqd_work_group_size() values\n");
goto err_free_dxil;
}
nir->info.cs.local_size[i] = conf->local_size[i];
nir->info.cs.workgroup_size[i] = conf->local_size[i];
}
memcpy(metadata->local_size, nir->info.cs.local_size,
memcpy(metadata->local_size, nir->info.cs.workgroup_size,
sizeof(metadata->local_size));
} else {
/* Make sure there's at least one thread that's set to run */
for (unsigned i = 0; i < ARRAY_SIZE(nir->info.cs.local_size); i++) {
if (nir->info.cs.local_size[i] == 0)
nir->info.cs.local_size[i] = 1;
for (unsigned i = 0; i < ARRAY_SIZE(nir->info.cs.workgroup_size); i++) {
if (nir->info.cs.workgroup_size[i] == 0)
nir->info.cs.workgroup_size[i] = 1;
}
}

6
src/microsoft/clc/clc_nir.c
View File

@ -73,9 +73,9 @@ lower_load_local_group_size(nir_builder *b, nir_intrinsic_instr *intr)
b->cursor = nir_after_instr(&intr->instr);
nir_const_value v[3] = {
nir_const_value_for_int(b->shader->info.cs.local_size[0], 32),
nir_const_value_for_int(b->shader->info.cs.local_size[1], 32),
nir_const_value_for_int(b->shader->info.cs.local_size[2], 32)
nir_const_value_for_int(b->shader->info.cs.workgroup_size[0], 32),
nir_const_value_for_int(b->shader->info.cs.workgroup_size[1], 32),
nir_const_value_for_int(b->shader->info.cs.workgroup_size[2], 32)
};
nir_ssa_def *size = nir_build_imm(b, 3, 32, v);
nir_ssa_def_rewrite_uses(&intr->dest.ssa, size);

6
src/microsoft/compiler/nir_to_dxil.c
View File

@ -1160,9 +1160,9 @@ static const struct dxil_mdnode *
emit_threads(struct ntd_context *ctx)
{
const nir_shader *s = ctx->shader;
const struct dxil_mdnode *threads_x = dxil_get_metadata_int32(&ctx->mod, MAX2(s->info.cs.local_size[0], 1));
const struct dxil_mdnode *threads_y = dxil_get_metadata_int32(&ctx->mod, MAX2(s->info.cs.local_size[1], 1));
const struct dxil_mdnode *threads_z = dxil_get_metadata_int32(&ctx->mod, MAX2(s->info.cs.local_size[2], 1));
const struct dxil_mdnode *threads_x = dxil_get_metadata_int32(&ctx->mod, MAX2(s->info.cs.workgroup_size[0], 1));
const struct dxil_mdnode *threads_y = dxil_get_metadata_int32(&ctx->mod, MAX2(s->info.cs.workgroup_size[1], 1));
const struct dxil_mdnode *threads_z = dxil_get_metadata_int32(&ctx->mod, MAX2(s->info.cs.workgroup_size[2], 1));
if (!threads_x || !threads_y || !threads_z)
return false;