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

nir: Rename WORK_GROUP (and similar) to WORKGROUP 

Be consistent with other usages in Vulkan and SPIR-V, and the recently
added workgroup_size field.

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-06-04 12:04:15 -07:00 committed by Marge Bot
parent a71a780598
commit c8a7bd0dc8
61 changed files with 176 additions and 176 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
src/amd/compiler/aco_instruction_selection.cpp
View File

@ -7818,7 +7818,7 @@ void visit_intrinsic(isel_context *ctx, nir_intrinsic_instr *instr)
case nir_intrinsic_scoped_barrier:
emit_scoped_barrier(ctx, instr);
break;
case nir_intrinsic_load_num_work_groups: {
case nir_intrinsic_load_num_workgroups: {
Temp dst = get_ssa_temp(ctx, &instr->dest.ssa);
bld.copy(Definition(dst), Operand(get_arg(ctx, ctx->args->ac.num_work_groups)));
emit_split_vector(ctx, dst, 3);
@ -7830,7 +7830,7 @@ void visit_intrinsic(isel_context *ctx, nir_intrinsic_instr *instr)
emit_split_vector(ctx, dst, 3);
break;
}
case nir_intrinsic_load_work_group_id: {
case nir_intrinsic_load_workgroup_id: {
Temp dst = get_ssa_temp(ctx, &instr->dest.ssa);
struct ac_arg *args = ctx->args->ac.workgroup_ids;
bld.pseudo(aco_opcode::p_create_vector, Definition(dst),

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

@ -530,13 +530,13 @@ void init_context(isel_context *ctx, nir_shader *shader)
ctx->ub_config.min_subgroup_size = ctx->options->key.cs.subgroup_size;
ctx->ub_config.max_subgroup_size = ctx->options->key.cs.subgroup_size;
}
ctx->ub_config.max_work_group_invocations = 2048;
ctx->ub_config.max_work_group_count[0] = 65535;
ctx->ub_config.max_work_group_count[1] = 65535;
ctx->ub_config.max_work_group_count[2] = 65535;
ctx->ub_config.max_work_group_size[0] = 2048;
ctx->ub_config.max_work_group_size[1] = 2048;
ctx->ub_config.max_work_group_size[2] = 2048;
ctx->ub_config.max_workgroup_invocations = 2048;
ctx->ub_config.max_workgroup_count[0] = 65535;
ctx->ub_config.max_workgroup_count[1] = 65535;
ctx->ub_config.max_workgroup_count[2] = 65535;
ctx->ub_config.max_workgroup_size[0] = 2048;
ctx->ub_config.max_workgroup_size[1] = 2048;
ctx->ub_config.max_workgroup_size[2] = 2048;
for (unsigned i = 0; i < MAX_VERTEX_ATTRIBS; i++) {
unsigned attrib_format = ctx->options->key.vs.vertex_attribute_formats[i];
unsigned dfmt = attrib_format & 0xf;
@ -719,8 +719,8 @@ void init_context(isel_context *ctx, nir_shader *shader)
RegType type = RegType::sgpr;
switch(intrinsic->intrinsic) {
case nir_intrinsic_load_push_constant:
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_workgroup_id:
case nir_intrinsic_load_num_workgroups:
case nir_intrinsic_load_subgroup_id:
case nir_intrinsic_load_num_subgroups:
case nir_intrinsic_load_first_vertex:

4
src/amd/llvm/ac_nir_to_llvm.c
View File

@ -3400,7 +3400,7 @@ static void visit_intrinsic(struct ac_nir_context *ctx, nir_intrinsic_instr *ins
case nir_intrinsic_load_subgroup_invocation:
result = ac_get_thread_id(&ctx->ac);
break;
case nir_intrinsic_load_work_group_id: {
case nir_intrinsic_load_workgroup_id: {
LLVMValueRef values[3];
for (int i = 0; i < 3; i++) {
@ -3516,7 +3516,7 @@ static void visit_intrinsic(struct ac_nir_context *ctx, nir_intrinsic_instr *ins
case nir_intrinsic_load_instance_id:
result = ctx->abi->instance_id;
break;
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_num_workgroups:
result = ac_get_arg(&ctx->ac, ctx->args->num_work_groups);
break;
case nir_intrinsic_load_local_invocation_index:

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

@ -13,7 +13,7 @@ build_buffer_fill_shader(struct radv_device *dev)
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 *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);
@ -43,7 +43,7 @@ build_buffer_copy_shader(struct radv_device *dev)
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 *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);

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

@ -52,7 +52,7 @@ build_nir_itob_compute_shader(struct radv_device *dev, bool is_3d)
output_img->data.binding = 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 *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);
@ -239,7 +239,7 @@ build_nir_btoi_compute_shader(struct radv_device *dev, bool is_3d)
output_img->data.binding = 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 *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);
@ -421,7 +421,7 @@ build_nir_btoi_r32g32b32_compute_shader(struct radv_device *dev)
output_img->data.binding = 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 *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);
@ -583,7 +583,7 @@ build_nir_itoi_compute_shader(struct radv_device *dev, bool is_3d, int samples)
output_img->data.binding = 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 *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);
@ -785,7 +785,7 @@ build_nir_itoi_r32g32b32_compute_shader(struct radv_device *dev)
output_img->data.binding = 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 *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);
@ -951,7 +951,7 @@ build_nir_cleari_compute_shader(struct radv_device *dev, bool is_3d, int samples
output_img->data.binding = 0;
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 *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);
@ -1117,7 +1117,7 @@ build_nir_cleari_r32g32b32_compute_shader(struct radv_device *dev)
output_img->data.binding = 0;
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 *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);

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

@ -1007,7 +1007,7 @@ build_clear_htile_mask_shader()
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 *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);

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

@ -50,7 +50,7 @@ build_copy_vrs_htile_shader(struct radv_device *device, struct radeon_surf *surf
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 *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);

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

@ -33,7 +33,7 @@ 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_work_group_id(b, 32), 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.cs.workgroup_size[0], b->shader->info.cs.workgroup_size[1],

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

@ -49,7 +49,7 @@ build_dcc_decompress_compute_shader(struct radv_device *dev)
output_img->data.binding = 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 *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);

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

@ -49,7 +49,7 @@ build_fmask_expand_compute_shader(struct radv_device *device, int samples)
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_work_group_id(&b, 32);
nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);

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

@ -79,7 +79,7 @@ build_resolve_compute_shader(struct radv_device *dev, bool is_integer, bool is_s
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_work_group_id(&b, 32);
nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);
@ -149,7 +149,7 @@ build_depth_stencil_resolve_compute_shader(struct radv_device *dev, 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_work_group_id(&b, 32);
nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);

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

@ -149,7 +149,7 @@ build_occlusion_query_shader(struct radv_device *device)
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_work_group_id(&b, 32);
nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);
@ -290,7 +290,7 @@ build_pipeline_statistics_query_shader(struct radv_device *device)
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_work_group_id(&b, 32);
nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);
@ -441,7 +441,7 @@ build_tfb_query_shader(struct radv_device *device)
/* Compute global ID. */
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 *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);
@ -571,7 +571,7 @@ build_timestamp_query_shader(struct radv_device *device)
/* Compute global ID. */
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 *wg_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *block_size =
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);

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

@ -158,16 +158,16 @@ gather_intrinsic_info(const nir_shader *nir, const nir_intrinsic_instr *instr,
case nir_intrinsic_load_instance_id:
info->vs.needs_instance_id = true;
break;
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_num_workgroups:
info->cs.uses_grid_size = true;
break;
case nir_intrinsic_load_local_invocation_id:
case nir_intrinsic_load_work_group_id: {
case nir_intrinsic_load_workgroup_id: {
unsigned mask = nir_ssa_def_components_read(&instr->dest.ssa);
while (mask) {
unsigned i = u_bit_scan(&mask);
if (instr->intrinsic == nir_intrinsic_load_work_group_id)
if (instr->intrinsic == nir_intrinsic_load_workgroup_id)
info->cs.uses_block_id[i] = true;
else
info->cs.uses_thread_id[i] = true;

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

@ -3026,7 +3026,7 @@ ntq_emit_intrinsic(struct v3d_compile *c, nir_intrinsic_instr *instr)
vir_emit_thrsw(c);
break;
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_num_workgroups:
for (int i = 0; i < 3; i++) {
ntq_store_dest(c, &instr->dest, i,
vir_uniform(c, QUNIFORM_NUM_WORK_GROUPS,
@ -3040,7 +3040,7 @@ ntq_emit_intrinsic(struct v3d_compile *c, nir_intrinsic_instr *instr)
vir_uniform_ui(c, 32 - c->local_invocation_index_bits)));
break;
case nir_intrinsic_load_work_group_id: {
case nir_intrinsic_load_workgroup_id: {
struct qreg x = vir_AND(c, c->cs_payload[0],
vir_uniform_ui(c, 0xffff));

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

@ -1452,8 +1452,8 @@ builtin_variable_generator::generate_cs_special_vars()
{
add_system_value(SYSTEM_VALUE_LOCAL_INVOCATION_ID, uvec3_t,
"gl_LocalInvocationID");
add_system_value(SYSTEM_VALUE_WORK_GROUP_ID, uvec3_t, "gl_WorkGroupID");
add_system_value(SYSTEM_VALUE_NUM_WORK_GROUPS, uvec3_t, "gl_NumWorkGroups");
add_system_value(SYSTEM_VALUE_WORKGROUP_ID, uvec3_t, "gl_WorkGroupID");
add_system_value(SYSTEM_VALUE_NUM_WORKGROUPS, uvec3_t, "gl_NumWorkGroups");
if (state->ARB_compute_variable_group_size_enable) {
add_system_value(SYSTEM_VALUE_WORKGROUP_SIZE,

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

@ -129,7 +129,7 @@ lower_cs_derived_visitor::find_sysvals()
if (!gl_WorkGroupID)
gl_WorkGroupID = add_system_value(
SYSTEM_VALUE_WORK_GROUP_ID, glsl_type::uvec3_type, "gl_WorkGroupID");
SYSTEM_VALUE_WORKGROUP_ID, glsl_type::uvec3_type, "gl_WorkGroupID");
if (!gl_LocalInvocationID)
gl_LocalInvocationID = add_system_value(
SYSTEM_VALUE_LOCAL_INVOCATION_ID, glsl_type::uvec3_type,

16
src/compiler/nir/nir.c
View File

@ -1975,10 +1975,10 @@ nir_intrinsic_from_system_value(gl_system_value val)
return nir_intrinsic_load_local_invocation_id;
case SYSTEM_VALUE_LOCAL_INVOCATION_INDEX:
return nir_intrinsic_load_local_invocation_index;
case SYSTEM_VALUE_WORK_GROUP_ID:
return nir_intrinsic_load_work_group_id;
case SYSTEM_VALUE_NUM_WORK_GROUPS:
return nir_intrinsic_load_num_work_groups;
case SYSTEM_VALUE_WORKGROUP_ID:
return nir_intrinsic_load_workgroup_id;
case SYSTEM_VALUE_NUM_WORKGROUPS:
return nir_intrinsic_load_num_workgroups;
case SYSTEM_VALUE_PRIMITIVE_ID:
return nir_intrinsic_load_primitive_id;
case SYSTEM_VALUE_TESS_COORD:
@ -2106,10 +2106,10 @@ nir_system_value_from_intrinsic(nir_intrinsic_op intrin)
return SYSTEM_VALUE_LOCAL_INVOCATION_ID;
case nir_intrinsic_load_local_invocation_index:
return SYSTEM_VALUE_LOCAL_INVOCATION_INDEX;
case nir_intrinsic_load_num_work_groups:
return SYSTEM_VALUE_NUM_WORK_GROUPS;
case nir_intrinsic_load_work_group_id:
return SYSTEM_VALUE_WORK_GROUP_ID;
case nir_intrinsic_load_num_workgroups:
return SYSTEM_VALUE_NUM_WORKGROUPS;
case nir_intrinsic_load_workgroup_id:
return SYSTEM_VALUE_WORKGROUP_ID;
case nir_intrinsic_load_primitive_id:
return SYSTEM_VALUE_PRIMITIVE_ID;
case nir_intrinsic_load_tess_coord:

10
src/compiler/nir/nir.h
View File

@ -3297,7 +3297,7 @@ typedef struct nir_shader_compiler_options {
bool lower_cs_local_index_from_id;
bool lower_cs_local_id_from_index;
/* Prevents lowering global_invocation_id to be in terms of work_group_id */
/* Prevents lowering global_invocation_id to be in terms of workgroup_id */
bool has_cs_global_id;
bool lower_device_index_to_zero;
@ -4674,7 +4674,7 @@ bool nir_lower_system_values(nir_shader *shader);
typedef struct nir_lower_compute_system_values_options {
bool has_base_global_invocation_id:1;
bool has_base_work_group_id:1;
bool has_base_workgroup_id:1;
bool shuffle_local_ids_for_quad_derivatives:1;
bool lower_local_invocation_index:1;
} nir_lower_compute_system_values_options;
@ -5256,9 +5256,9 @@ nir_variable_is_in_block(const nir_variable *var)
typedef struct nir_unsigned_upper_bound_config {
unsigned min_subgroup_size;
unsigned max_subgroup_size;
unsigned max_work_group_invocations;
unsigned max_work_group_count[3];
unsigned max_work_group_size[3];
unsigned max_workgroup_invocations;
unsigned max_workgroup_count[3];
unsigned max_workgroup_size[3];
uint32_t vertex_attrib_max[32];
} nir_unsigned_upper_bound_config;

4
src/compiler/nir/nir_divergence_analysis.c
View File

@ -101,7 +101,7 @@ visit_intrinsic(nir_shader *shader, nir_intrinsic_instr *instr)
case nir_intrinsic_vote_ieq:
case nir_intrinsic_load_push_constant:
case nir_intrinsic_load_work_dim:
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_num_workgroups:
case nir_intrinsic_load_workgroup_size:
case nir_intrinsic_load_subgroup_id:
case nir_intrinsic_load_num_subgroups:
@ -236,7 +236,7 @@ visit_intrinsic(nir_shader *shader, nir_intrinsic_instr *instr)
assert(stage == MESA_SHADER_TESS_CTRL);
break;
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_workgroup_id:
assert(stage == MESA_SHADER_COMPUTE);
is_divergent |= (options & nir_divergence_multiple_workgroup_per_compute_subgroup);
break;

4
src/compiler/nir/nir_gather_info.c
View File

@ -623,8 +623,8 @@ gather_intrinsic_info(nir_intrinsic_instr *instr, nir_shader *shader,
case nir_intrinsic_load_global_invocation_id:
case nir_intrinsic_load_base_global_invocation_id:
case nir_intrinsic_load_global_invocation_index:
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_workgroup_id:
case nir_intrinsic_load_num_workgroups:
case nir_intrinsic_load_workgroup_size:
case nir_intrinsic_load_work_dim:
case nir_intrinsic_load_user_data_amd:

12
src/compiler/nir/nir_intrinsics.py
View File

@ -685,11 +685,11 @@ system_value("local_invocation_id", 3)
system_value("local_invocation_index", 1)
# zero_base indicates it starts from 0 for the current dispatch
# non-zero_base indicates the base is included
system_value("work_group_id", 3, bit_sizes=[32, 64])
system_value("work_group_id_zero_base", 3)
system_value("base_work_group_id", 3, bit_sizes=[32, 64])
system_value("workgroup_id", 3, bit_sizes=[32, 64])
system_value("workgroup_id_zero_base", 3)
system_value("base_workgroup_id", 3, bit_sizes=[32, 64])
system_value("user_clip_plane", 4, indices=[UCP_ID])
system_value("num_work_groups", 3, bit_sizes=[32, 64])
system_value("num_workgroups", 3, bit_sizes=[32, 64])
system_value("helper_invocation", 1, bit_sizes=[1, 32])
system_value("layer_id", 1)
system_value("view_index", 1)
@ -704,8 +704,8 @@ system_value("num_subgroups", 1)
system_value("subgroup_id", 1)
system_value("workgroup_size", 3)
# note: the definition of global_invocation_id_zero_base is based on
# (work_group_id * workgroup_size) + local_invocation_id.
# it is *not* based on work_group_id_zero_base, meaning the work group
# (workgroup_id * workgroup_size) + local_invocation_id.
# it is *not* based on workgroup_id_zero_base, meaning the work group
# base is already accounted for, and the global base is additive on top of that
system_value("global_invocation_id", 3, bit_sizes=[32, 64])
system_value("global_invocation_id_zero_base", 3, bit_sizes=[32, 64])

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

@ -55,9 +55,9 @@ static nir_ssa_def*
build_global_group_size(nir_builder *b, unsigned bit_size)
{
nir_ssa_def *group_size = nir_load_workgroup_size(b);
nir_ssa_def *num_work_groups = nir_load_num_work_groups(b, bit_size);
nir_ssa_def *num_workgroups = nir_load_num_workgroups(b, bit_size);
return nir_imul(b, nir_u2u(b, group_size, bit_size),
num_work_groups);
num_workgroups);
}
static bool
@ -443,10 +443,10 @@ lower_compute_system_value_instr(nir_builder *b,
}
case nir_intrinsic_load_global_invocation_id_zero_base: {
if ((options && options->has_base_work_group_id) ||
if ((options && options->has_base_workgroup_id) ||
!b->shader->options->has_cs_global_id) {
nir_ssa_def *group_size = nir_load_workgroup_size(b);
nir_ssa_def *group_id = nir_load_work_group_id(b, bit_size);
nir_ssa_def *group_id = nir_load_workgroup_id(b, bit_size);
nir_ssa_def *local_id = nir_load_local_invocation_id(b);
return nir_iadd(b, nir_imul(b, group_id,
@ -461,7 +461,7 @@ lower_compute_system_value_instr(nir_builder *b,
if (options && options->has_base_global_invocation_id)
return nir_iadd(b, nir_load_global_invocation_id_zero_base(b, bit_size),
nir_load_base_global_invocation_id(b, bit_size));
else if ((options && options->has_base_work_group_id) ||
else if ((options && options->has_base_workgroup_id) ||
!b->shader->options->has_cs_global_id)
return nir_load_global_invocation_id_zero_base(b, bit_size);
else
@ -485,10 +485,10 @@ lower_compute_system_value_instr(nir_builder *b,
return index;
}
case nir_intrinsic_load_work_group_id: {
if (options && options->has_base_work_group_id)
return nir_iadd(b, nir_u2u(b, nir_load_work_group_id_zero_base(b), bit_size),
nir_load_base_work_group_id(b, bit_size));
case nir_intrinsic_load_workgroup_id: {
if (options && options->has_base_workgroup_id)
return nir_iadd(b, nir_u2u(b, nir_load_workgroup_id_zero_base(b), bit_size),
nir_load_base_workgroup_id(b, bit_size));
else
return NULL;
}

4
src/compiler/nir/nir_opt_peephole_select.c
View File

@ -106,8 +106,8 @@ block_check_for_allowed_instrs(nir_block *block, unsigned *count,
case nir_intrinsic_load_base_instance:
case nir_intrinsic_load_instance_id:
case nir_intrinsic_load_draw_id:
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_num_workgroups:
case nir_intrinsic_load_workgroup_id:
case nir_intrinsic_load_local_invocation_id:
case nir_intrinsic_load_local_invocation_index:
case nir_intrinsic_load_subgroup_id:

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

@ -1256,9 +1256,9 @@ lookup_input(nir_shader *shader, unsigned driver_location)
static const nir_unsigned_upper_bound_config default_ub_config = {
.min_subgroup_size = 1u,
.max_subgroup_size = UINT16_MAX,
.max_work_group_invocations = UINT16_MAX,
.max_work_group_count = {UINT16_MAX, UINT16_MAX, UINT16_MAX},
.max_work_group_size = {UINT16_MAX, UINT16_MAX, UINT16_MAX},
.max_workgroup_invocations = UINT16_MAX,
.max_workgroup_count = {UINT16_MAX, UINT16_MAX, UINT16_MAX},
.max_workgroup_size = {UINT16_MAX, UINT16_MAX, UINT16_MAX},
.vertex_attrib_max = {
UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX,
UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX,
@ -1294,7 +1294,7 @@ nir_unsigned_upper_bound(nir_shader *shader, struct hash_table *range_ht,
case nir_intrinsic_load_local_invocation_index:
if (shader->info.stage != MESA_SHADER_COMPUTE ||
shader->info.cs.workgroup_size_variable) {
res = config->max_work_group_invocations - 1;
res = config->max_workgroup_invocations - 1;
} else {
res = (shader->info.cs.workgroup_size[0] *
shader->info.cs.workgroup_size[1] *
@ -1303,23 +1303,23 @@ nir_unsigned_upper_bound(nir_shader *shader, struct hash_table *range_ht,
break;
case nir_intrinsic_load_local_invocation_id:
if (shader->info.cs.workgroup_size_variable)
res = config->max_work_group_size[scalar.comp] - 1u;
res = config->max_workgroup_size[scalar.comp] - 1u;
else
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;
case nir_intrinsic_load_workgroup_id:
res = config->max_workgroup_count[scalar.comp] - 1u;
break;
case nir_intrinsic_load_num_work_groups:
res = config->max_work_group_count[scalar.comp];
case nir_intrinsic_load_num_workgroups:
res = config->max_workgroup_count[scalar.comp];
break;
case nir_intrinsic_load_global_invocation_id:
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;
res = mul_clamp(config->max_workgroup_size[scalar.comp],
config->max_workgroup_count[scalar.comp]) - 1u;
} else {
res = (shader->info.cs.workgroup_size[scalar.comp] *
config->max_work_group_count[scalar.comp]) - 1u;
config->max_workgroup_count[scalar.comp]) - 1u;
}
break;
case nir_intrinsic_load_invocation_id:
@ -1338,13 +1338,13 @@ nir_unsigned_upper_bound(nir_shader *shader, struct hash_table *range_ht,
break;
case nir_intrinsic_load_subgroup_id:
case nir_intrinsic_load_num_subgroups: {
uint32_t work_group_size = config->max_work_group_invocations;
uint32_t workgroup_size = config->max_workgroup_invocations;
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];
workgroup_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);
res = DIV_ROUND_UP(workgroup_size, config->min_subgroup_size);
if (intrin->intrinsic == nir_intrinsic_load_subgroup_id)
res--;
break;
@ -1391,7 +1391,7 @@ nir_unsigned_upper_bound(nir_shader *shader, struct hash_table *range_ht,
case nir_intrinsic_load_tess_rel_patch_id_amd:
case nir_intrinsic_load_tcs_num_patches_amd:
/* Very generous maximum: TCS/TES executed by largest possible workgroup */
res = config->max_work_group_invocations / MAX2(shader->info.tess.tcs_vertices_out, 1u);
res = config->max_workgroup_invocations / MAX2(shader->info.tess.tcs_vertices_out, 1u);
break;
default:
break;

4
src/compiler/shader_enums.c
View File

@ -275,8 +275,8 @@ gl_system_value_name(gl_system_value sysval)
ENUM(SYSTEM_VALUE_GLOBAL_INVOCATION_ID),
ENUM(SYSTEM_VALUE_BASE_GLOBAL_INVOCATION_ID),
ENUM(SYSTEM_VALUE_GLOBAL_INVOCATION_INDEX),
ENUM(SYSTEM_VALUE_WORK_GROUP_ID),
ENUM(SYSTEM_VALUE_NUM_WORK_GROUPS),
ENUM(SYSTEM_VALUE_WORKGROUP_ID),
ENUM(SYSTEM_VALUE_NUM_WORKGROUPS),
ENUM(SYSTEM_VALUE_WORKGROUP_SIZE),
ENUM(SYSTEM_VALUE_GLOBAL_GROUP_SIZE),
ENUM(SYSTEM_VALUE_USER_DATA_AMD),

4
src/compiler/shader_enums.h
View File

@ -710,8 +710,8 @@ typedef enum
SYSTEM_VALUE_GLOBAL_INVOCATION_ID,
SYSTEM_VALUE_BASE_GLOBAL_INVOCATION_ID,
SYSTEM_VALUE_GLOBAL_INVOCATION_INDEX,
SYSTEM_VALUE_WORK_GROUP_ID,
SYSTEM_VALUE_NUM_WORK_GROUPS,
SYSTEM_VALUE_WORKGROUP_ID,
SYSTEM_VALUE_NUM_WORKGROUPS,
SYSTEM_VALUE_WORKGROUP_SIZE,
SYSTEM_VALUE_GLOBAL_GROUP_SIZE,
SYSTEM_VALUE_WORK_DIM,

4
src/compiler/spirv/vtn_variables.c
View File

@ -914,7 +914,7 @@ vtn_get_builtin_location(struct vtn_builder *b,
set_mode_system_value(b, mode);
break;
case SpvBuiltInNumWorkgroups:
*location = SYSTEM_VALUE_NUM_WORK_GROUPS;
*location = SYSTEM_VALUE_NUM_WORKGROUPS;
set_mode_system_value(b, mode);
break;
case SpvBuiltInWorkgroupSize:
@ -923,7 +923,7 @@ vtn_get_builtin_location(struct vtn_builder *b,
set_mode_system_value(b, mode);
break;
case SpvBuiltInWorkgroupId:
*location = SYSTEM_VALUE_WORK_GROUP_ID;
*location = SYSTEM_VALUE_WORKGROUP_ID;
set_mode_system_value(b, mode);
break;
case SpvBuiltInLocalInvocationId:

2
src/freedreno/computerator/a6xx.c
View File

@ -156,7 +156,7 @@ cs_program_emit(struct fd_ringbuffer *ring, struct kernel *kernel)
uint32_t local_invocation_id, work_group_id;
local_invocation_id =
ir3_find_sysval_regid(v, SYSTEM_VALUE_LOCAL_INVOCATION_ID);
work_group_id = ir3_find_sysval_regid(v, SYSTEM_VALUE_WORK_GROUP_ID);
work_group_id = ir3_find_sysval_regid(v, SYSTEM_VALUE_WORKGROUP_ID);
OUT_PKT4(ring, REG_A6XX_HLSQ_CS_CNTL_0, 2);
OUT_RING(ring, A6XX_HLSQ_CS_CNTL_0_WGIDCONSTID(work_group_id) |

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

@ -2023,21 +2023,21 @@ emit_intrinsic(struct ir3_context *ctx, nir_intrinsic_instr *intr)
}
ir3_split_dest(b, dst, ctx->local_invocation_id, 0, 3);
break;
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_work_group_id_zero_base:
case nir_intrinsic_load_workgroup_id:
case nir_intrinsic_load_workgroup_id_zero_base:
if (!ctx->work_group_id) {
ctx->work_group_id =
create_sysval_input(ctx, SYSTEM_VALUE_WORK_GROUP_ID, 0x7);
create_sysval_input(ctx, SYSTEM_VALUE_WORKGROUP_ID, 0x7);
ctx->work_group_id->regs[0]->flags |= IR3_REG_SHARED;
}
ir3_split_dest(b, dst, ctx->work_group_id, 0, 3);
break;
case nir_intrinsic_load_base_work_group_id:
case nir_intrinsic_load_base_workgroup_id:
for (int i = 0; i < dest_components; i++) {
dst[i] = create_driver_param(ctx, IR3_DP_BASE_GROUP_X + i);
}
break;
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_num_workgroups:
for (int i = 0; i < dest_components; i++) {
dst[i] = create_driver_param(ctx, IR3_DP_NUM_WORK_GROUPS_X + i);
}

4
src/freedreno/ir3/ir3_nir.c
View File

@ -676,7 +676,7 @@ ir3_nir_scan_driver_consts(nir_shader *shader,
layout->num_driver_params =
MAX2(layout->num_driver_params, IR3_DP_UCP0_X + (idx + 1) * 4);
break;
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_num_workgroups:
layout->num_driver_params =
MAX2(layout->num_driver_params, IR3_DP_NUM_WORK_GROUPS_Z + 1);
break;
@ -684,7 +684,7 @@ ir3_nir_scan_driver_consts(nir_shader *shader,
layout->num_driver_params =
MAX2(layout->num_driver_params, IR3_DP_LOCAL_GROUP_SIZE_Z + 1);
break;
case nir_intrinsic_load_base_work_group_id:
case nir_intrinsic_load_base_workgroup_id:
layout->num_driver_params =
MAX2(layout->num_driver_params, IR3_DP_BASE_GROUP_Z + 1);
break;

2
src/freedreno/ir3/ir3_parser.y
View File

@ -671,7 +671,7 @@ wgid_header: T_A_WGID '(' T_REGISTER ')' {
assert(($3 & 0x1) == 0); /* half-reg not allowed */
unsigned reg = $3 >> 1;
assert(reg >= regid(48, 0)); /* must be a high reg */
add_sysval(reg, 0x7, SYSTEM_VALUE_WORK_GROUP_ID);
add_sysval(reg, 0x7, SYSTEM_VALUE_WORKGROUP_ID);
}
numwg_header: T_A_NUMWG '(' T_CONSTANT ')' {

2
src/freedreno/vulkan/tu_pipeline.c
View File

@ -578,7 +578,7 @@ tu6_emit_cs_config(struct tu_cs *cs, const struct tu_shader *shader,
uint32_t local_invocation_id =
ir3_find_sysval_regid(v, SYSTEM_VALUE_LOCAL_INVOCATION_ID);
uint32_t work_group_id =
ir3_find_sysval_regid(v, SYSTEM_VALUE_WORK_GROUP_ID);
ir3_find_sysval_regid(v, SYSTEM_VALUE_WORKGROUP_ID);
enum a6xx_threadsize thrsz = v->info.double_threadsize ? THREAD128 : THREAD64;
tu_cs_emit_pkt4(cs, REG_A6XX_HLSQ_CS_CNTL_0, 2);

2
src/freedreno/vulkan/tu_shader.c
View File

@ -84,7 +84,7 @@ tu_spirv_to_nir(struct tu_device *dev,
};
const struct nir_lower_compute_system_values_options compute_sysval_options = {
.has_base_work_group_id = true,
.has_base_workgroup_id = true,
};
const nir_shader_compiler_options *nir_options =

4
src/gallium/auxiliary/gallivm/lp_bld_nir.c
View File

@ -1690,9 +1690,9 @@ static void visit_intrinsic(struct lp_build_nir_context *bld_base,
case nir_intrinsic_load_base_instance:
case nir_intrinsic_load_base_vertex:
case nir_intrinsic_load_first_vertex:
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_workgroup_id:
case nir_intrinsic_load_local_invocation_id:
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_num_workgroups:
case nir_intrinsic_load_invocation_id:
case nir_intrinsic_load_front_face:
case nir_intrinsic_load_draw_id:

4
src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c
View File

@ -1514,7 +1514,7 @@ static void emit_sysval_intrin(struct lp_build_nir_context *bld_base,
case nir_intrinsic_load_primitive_id:
result[0] = bld->system_values.prim_id;
break;
case nir_intrinsic_load_work_group_id: {
case nir_intrinsic_load_workgroup_id: {
LLVMValueRef tmp[3];
for (unsigned i = 0; i < 3; i++) {
tmp[i] = LLVMBuildExtractElement(gallivm->builder, bld->system_values.block_id, lp_build_const_int32(gallivm, i), "");
@ -1528,7 +1528,7 @@ static void emit_sysval_intrin(struct lp_build_nir_context *bld_base,
for (unsigned i = 0; i < 3; i++)
result[i] = LLVMBuildExtractValue(gallivm->builder, bld->system_values.thread_id, i, "");
break;
case nir_intrinsic_load_num_work_groups: {
case nir_intrinsic_load_num_workgroups: {
LLVMValueRef tmp[3];
for (unsigned i = 0; i < 3; i++) {
tmp[i] = LLVMBuildExtractElement(gallivm->builder, bld->system_values.grid_size, lp_build_const_int32(gallivm, i), "");

4
src/gallium/auxiliary/nir/nir_to_tgsi.c
View File

@ -1650,8 +1650,8 @@ ntt_emit_intrinsic(struct ntt_compile *c, nir_intrinsic_instr *instr)
case nir_intrinsic_load_tess_level_outer:
case nir_intrinsic_load_tess_level_inner:
case nir_intrinsic_load_local_invocation_id:
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_workgroup_id:
case nir_intrinsic_load_num_workgroups:
case nir_intrinsic_load_workgroup_size:
case nir_intrinsic_load_subgroup_size:
case nir_intrinsic_load_subgroup_invocation:

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

@ -220,7 +220,7 @@ static void scan_instruction(const struct nir_shader *nir,
case nir_intrinsic_load_invocation_id:
info->uses_invocationid = true;
break;
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_num_workgroups:
info->uses_grid_size = true;
break;
case nir_intrinsic_load_workgroup_size:
@ -229,12 +229,12 @@ static void scan_instruction(const struct nir_shader *nir,
info->uses_block_size = true;
break;
case nir_intrinsic_load_local_invocation_id:
case nir_intrinsic_load_work_group_id: {
case nir_intrinsic_load_workgroup_id: {
unsigned mask = nir_ssa_def_components_read(&intr->dest.ssa);
while (mask) {
unsigned i = u_bit_scan(&mask);
if (intr->intrinsic == nir_intrinsic_load_work_group_id)
if (intr->intrinsic == nir_intrinsic_load_workgroup_id)
info->uses_block_id[i] = true;
else
info->uses_thread_id[i] = true;

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

@ -620,7 +620,7 @@ ttn_src_for_file_and_index(struct ttn_compile *c, unsigned file, unsigned index,
load = nir_load_local_invocation_id(b);
break;
case TGSI_SEMANTIC_BLOCK_ID:
load = nir_load_work_group_id(b, 32);
load = nir_load_workgroup_id(b, 32);
break;
case TGSI_SEMANTIC_BLOCK_SIZE:
load = nir_load_workgroup_size(b);

4
src/gallium/auxiliary/tgsi/tgsi_from_mesa.c
View File

@ -265,9 +265,9 @@ tgsi_get_sysval_semantic(unsigned sysval)
/* Compute shader */
case SYSTEM_VALUE_LOCAL_INVOCATION_ID:
return TGSI_SEMANTIC_THREAD_ID;
case SYSTEM_VALUE_WORK_GROUP_ID:
case SYSTEM_VALUE_WORKGROUP_ID:
return TGSI_SEMANTIC_BLOCK_ID;
case SYSTEM_VALUE_NUM_WORK_GROUPS:
case SYSTEM_VALUE_NUM_WORKGROUPS:
return TGSI_SEMANTIC_GRID_SIZE;
case SYSTEM_VALUE_WORKGROUP_SIZE:
return TGSI_SEMANTIC_BLOCK_SIZE;

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

@ -91,7 +91,7 @@ cs_program_emit(struct fd_ringbuffer *ring, struct ir3_shader_variant *v)
uint32_t local_invocation_id, work_group_id;
local_invocation_id =
ir3_find_sysval_regid(v, SYSTEM_VALUE_LOCAL_INVOCATION_ID);
work_group_id = ir3_find_sysval_regid(v, SYSTEM_VALUE_WORK_GROUP_ID);
work_group_id = ir3_find_sysval_regid(v, SYSTEM_VALUE_WORKGROUP_ID);
OUT_PKT4(ring, REG_A5XX_HLSQ_CS_CNTL_0, 2);
OUT_RING(ring, A5XX_HLSQ_CS_CNTL_0_WGIDCONSTID(work_group_id) |

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

@ -78,7 +78,7 @@ cs_program_emit(struct fd_context *ctx, struct fd_ringbuffer *ring,
uint32_t local_invocation_id, work_group_id;
local_invocation_id =
ir3_find_sysval_regid(v, SYSTEM_VALUE_LOCAL_INVOCATION_ID);
work_group_id = ir3_find_sysval_regid(v, SYSTEM_VALUE_WORK_GROUP_ID);
work_group_id = ir3_find_sysval_regid(v, SYSTEM_VALUE_WORKGROUP_ID);
OUT_PKT4(ring, REG_A6XX_HLSQ_CS_CNTL_0, 2);
OUT_RING(ring, A6XX_HLSQ_CS_CNTL_0_WGIDCONSTID(work_group_id) |

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

@ -852,7 +852,7 @@ iris_setup_binding_table(const struct intel_device_info *devinfo,
nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr);
switch (intrin->intrinsic) {
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_num_workgroups:
bt->used_mask[IRIS_SURFACE_GROUP_CS_WORK_GROUPS] = 1;
break;

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

@ -1570,7 +1570,7 @@ Converter::convert(nir_intrinsic_op intr)
return SV_NTID;
case nir_intrinsic_load_local_invocation_id:
return SV_TID;
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_num_workgroups:
return SV_NCTAID;
case nir_intrinsic_load_patch_vertices_in:
return SV_VERTEX_COUNT;
@ -1602,7 +1602,7 @@ Converter::convert(nir_intrinsic_op intr)
return SV_TESS_OUTER;
case nir_intrinsic_load_vertex_id:
return SV_VERTEX_ID;
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_workgroup_id:
return SV_CTAID;
case nir_intrinsic_load_work_dim:
return SV_WORK_DIM;
@ -1845,7 +1845,7 @@ Converter::visit(nir_intrinsic_instr *insn)
case nir_intrinsic_load_invocation_id:
case nir_intrinsic_load_workgroup_size:
case nir_intrinsic_load_local_invocation_id:
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_num_workgroups:
case nir_intrinsic_load_patch_vertices_in:
case nir_intrinsic_load_primitive_id:
case nir_intrinsic_load_sample_id:
@ -1861,7 +1861,7 @@ Converter::visit(nir_intrinsic_instr *insn)
case nir_intrinsic_load_tess_level_inner:
case nir_intrinsic_load_tess_level_outer:
case nir_intrinsic_load_vertex_id:
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_workgroup_id:
case nir_intrinsic_load_work_dim: {
const DataType dType = getDType(insn);
SVSemantic sv = convert(op);

8
src/gallium/drivers/r600/sfn/sfn_shader_compute.cpp
View File

@ -69,10 +69,10 @@ bool ComputeShaderFromNir::emit_intrinsic_instruction_override(nir_intrinsic_ins
switch (instr->intrinsic) {
case nir_intrinsic_load_local_invocation_id:
return emit_load_3vec(instr, m_local_invocation_id);
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_workgroup_id:
return emit_load_3vec(instr, m_workgroup_id);
case nir_intrinsic_load_num_work_groups:
return emit_load_num_work_groups(instr);
case nir_intrinsic_load_num_workgroups:
return emit_load_num_workgroups(instr);
default:
return false;
}
@ -86,7 +86,7 @@ bool ComputeShaderFromNir::emit_load_3vec(nir_intrinsic_instr* instr,
return true;
}
bool ComputeShaderFromNir::emit_load_num_work_groups(nir_intrinsic_instr* instr)
bool ComputeShaderFromNir::emit_load_num_workgroups(nir_intrinsic_instr* instr)
{
PValue a_zero = get_temp_register(1);
emit_instruction(new AluInstruction(op1_mov, a_zero, Value::zero, EmitInstruction::last_write));

2
src/gallium/drivers/r600/sfn/sfn_shader_compute.h
View File

@ -50,7 +50,7 @@ private:
void do_finalize() override;
bool emit_load_3vec(nir_intrinsic_instr* instr, const std::array<PValue,3>& src);
bool emit_load_num_work_groups(nir_intrinsic_instr* instr);
bool emit_load_num_workgroups(nir_intrinsic_instr* instr);
int m_reserved_registers;
std::array<PValue,3> m_workgroup_id;

6
src/gallium/drivers/radeonsi/si_shader_nir.c
View File

@ -301,12 +301,12 @@ static void scan_instruction(const struct nir_shader *nir, struct si_shader_info
info->uses_indirect_descriptor = true;
break;
case nir_intrinsic_load_local_invocation_id:
case nir_intrinsic_load_work_group_id: {
case nir_intrinsic_load_workgroup_id: {
unsigned mask = nir_ssa_def_components_read(&intr->dest.ssa);
while (mask) {
unsigned i = u_bit_scan(&mask);
if (intr->intrinsic == nir_intrinsic_load_work_group_id)
if (intr->intrinsic == nir_intrinsic_load_workgroup_id)
info->uses_block_id[i] = true;
else
info->uses_thread_id[i] = true;
@ -431,7 +431,7 @@ void si_nir_scan_shader(const struct nir_shader *nir, struct si_shader_info *inf
info->uses_base_vertex = BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_BASE_VERTEX);
info->uses_base_instance = BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_BASE_INSTANCE);
info->uses_invocationid = BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_INVOCATION_ID);
info->uses_grid_size = BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_NUM_WORK_GROUPS);
info->uses_grid_size = BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_NUM_WORKGROUPS);
info->uses_subgroup_info = BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_LOCAL_INVOCATION_INDEX) ||
BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_SUBGROUP_ID) ||
BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_NUM_SUBGROUPS);

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

@ -42,7 +42,7 @@ 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_work_group_id(b, 32), 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_load_workgroup_size(b), mask);
return nir_iadd(b, nir_imul(b, block_ids, block_size), local_ids);
}

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

@ -2765,11 +2765,11 @@ emit_intrinsic(struct ntv_context *ctx, nir_intrinsic_instr *intr)
emit_image_intrinsic(ctx, intr);
break;
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_workgroup_id:
emit_load_vec_input(ctx, intr, &ctx->workgroup_id_var, "gl_WorkGroupID", SpvBuiltInWorkgroupId, nir_type_uint);
break;
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_num_workgroups:
emit_load_vec_input(ctx, intr, &ctx->num_workgroups_var, "gl_NumWorkGroups", SpvBuiltInNumWorkgroups, nir_type_uint);
break;

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

@ -191,10 +191,10 @@ emit_system_values_block(nir_block *block, fs_visitor *v)
*reg = *v->emit_samplemaskin_setup();
break;
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_workgroup_id:
assert(v->stage == MESA_SHADER_COMPUTE ||
v->stage == MESA_SHADER_KERNEL);
reg = &v->nir_system_values[SYSTEM_VALUE_WORK_GROUP_ID];
reg = &v->nir_system_values[SYSTEM_VALUE_WORKGROUP_ID];
if (reg->file == BAD_FILE)
*reg = *v->emit_cs_work_group_id_setup();
break;
@ -3691,7 +3691,7 @@ fs_visitor::nir_emit_cs_intrinsic(const fs_builder &bld,
break;
case nir_intrinsic_load_local_invocation_id:
case nir_intrinsic_load_work_group_id: {
case nir_intrinsic_load_workgroup_id: {
gl_system_value sv = nir_system_value_from_intrinsic(instr->intrinsic);
fs_reg val = nir_system_values[sv];
assert(val.file != BAD_FILE);
@ -3701,7 +3701,7 @@ fs_visitor::nir_emit_cs_intrinsic(const fs_builder &bld,
break;
}
case nir_intrinsic_load_num_work_groups: {
case nir_intrinsic_load_num_workgroups: {
assert(nir_dest_bit_size(instr->dest) == 32);
const unsigned surface =
cs_prog_data->binding_table.work_groups_start;

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

@ -54,8 +54,8 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state,
nir_ssa_def *sysval;
switch (intrinsic->intrinsic) {
case nir_intrinsic_load_workgroup_size:
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_workgroup_id:
case nir_intrinsic_load_num_workgroups:
/* Convert this to 32-bit if it's not */
if (intrinsic->dest.ssa.bit_size == 64) {
intrinsic->dest.ssa.bit_size = 32;

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

@ -438,7 +438,7 @@ brw_nir_create_raygen_trampoline(const struct brw_compiler *compiler,
nir_ssa_def *local_shift =
nir_u2u32(&b, load_trampoline_param(&b, local_group_size_log2, 3, 8));
nir_ssa_def *global_id = nir_load_work_group_id(&b, 32);
nir_ssa_def *global_id = nir_load_workgroup_id(&b, 32);
nir_ssa_def *simd_channel = nir_load_subgroup_invocation(&b);
nir_ssa_def *local_x =
nir_ubfe(&b, simd_channel, nir_imm_int(&b, 0),

2
src/intel/vulkan/anv_nir_add_base_work_group_id.c
View File

@ -44,7 +44,7 @@ anv_nir_add_base_work_group_id(nir_shader *shader)
continue;
nir_intrinsic_instr *load_id = nir_instr_as_intrinsic(instr);
if (load_id->intrinsic != nir_intrinsic_load_work_group_id)
if (load_id->intrinsic != nir_intrinsic_load_workgroup_id)
continue;
b.cursor = nir_after_instr(&load_id->instr);

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

@ -351,10 +351,10 @@ clc_lower_64bit_semantics(nir_shader *nir)
case nir_intrinsic_load_global_invocation_id_zero_base:
case nir_intrinsic_load_base_global_invocation_id:
case nir_intrinsic_load_local_invocation_id:
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_work_group_id_zero_base:
case nir_intrinsic_load_base_work_group_id:
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_workgroup_id:
case nir_intrinsic_load_workgroup_id_zero_base:
case nir_intrinsic_load_base_workgroup_id:
case nir_intrinsic_load_num_workgroups:
break;
default:
continue;
@ -1322,7 +1322,7 @@ clc_to_dxil(struct clc_context *ctx,
nir_lower_compute_system_values_options compute_options = {
.has_base_global_invocation_id = (conf && conf->support_global_work_id_offsets),
.has_base_work_group_id = (conf && conf->support_work_group_id_offsets),
.has_base_workgroup_id = (conf && conf->support_workgroup_id_offsets),
};
NIR_PASS_V(nir, nir_lower_compute_system_values, &compute_options);

2
src/microsoft/clc/clc_compiler.h
View File

@ -233,7 +233,7 @@ struct clc_runtime_kernel_conf {
struct clc_runtime_arg_info *args;
unsigned lower_bit_size;
unsigned support_global_work_id_offsets;
unsigned support_work_group_id_offsets;
unsigned support_workgroup_id_offsets;
};
struct clc_dxil_object *

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

@ -84,8 +84,8 @@ lower_load_local_group_size(nir_builder *b, nir_intrinsic_instr *intr)
}
static bool
lower_load_num_work_groups(nir_builder *b, nir_intrinsic_instr *intr,
nir_variable *var)
lower_load_num_workgroups(nir_builder *b, nir_intrinsic_instr *intr,
nir_variable *var)
{
b->cursor = nir_after_instr(&intr->instr);
@ -102,7 +102,7 @@ lower_load_num_work_groups(nir_builder *b, nir_intrinsic_instr *intr,
}
static bool
lower_load_base_work_group_id(nir_builder *b, nir_intrinsic_instr *intr,
lower_load_base_workgroup_id(nir_builder *b, nir_intrinsic_instr *intr,
nir_variable *var)
{
b->cursor = nir_after_instr(&intr->instr);
@ -149,11 +149,11 @@ clc_nir_lower_system_values(nir_shader *nir, nir_variable *var)
case nir_intrinsic_load_workgroup_size:
lower_load_local_group_size(&b, intr);
break;
case nir_intrinsic_load_num_work_groups:
lower_load_num_work_groups(&b, intr, var);
case nir_intrinsic_load_num_workgroups:
lower_load_num_workgroups(&b, intr, var);
break;
case nir_intrinsic_load_base_work_group_id:
lower_load_base_work_group_id(&b, intr, var);
case nir_intrinsic_load_base_workgroup_id:
lower_load_base_workgroup_id(&b, intr, var);
break;
default: break;
}

4
src/microsoft/clc/compute_test.cpp
View File

@ -467,7 +467,7 @@ ComputeTest::run_shader_with_raw_args(Shader shader,
compile_args.work_props.global_offset_x != 0 ||
compile_args.work_props.global_offset_y != 0 ||
compile_args.work_props.global_offset_z != 0;
conf.support_work_group_id_offsets =
conf.support_workgroup_id_offsets =
compile_args.work_props.group_id_offset_x != 0 ||
compile_args.work_props.group_id_offset_y != 0 ||
compile_args.work_props.group_id_offset_z != 0;
@ -493,7 +493,7 @@ ComputeTest::run_shader_with_raw_args(Shader shader,
std::vector<uint8_t> argsbuf(dxil->metadata.kernel_inputs_buf_size);
std::vector<ComPtr<ID3D12Resource>> argres(shader.dxil->kernel->num_args);
clc_work_properties_data work_props = compile_args.work_props;
if (!conf.support_work_group_id_offsets) {
if (!conf.support_workgroup_id_offsets) {
work_props.group_count_total_x = compile_args.x / conf.local_size[0];
work_props.group_count_total_y = compile_args.y / conf.local_size[1];
work_props.group_count_total_z = compile_args.z / conf.local_size[2];

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

@ -2264,7 +2264,7 @@ emit_load_local_invocation_id(struct ntd_context *ctx,
}
static bool
emit_load_local_work_group_id(struct ntd_context *ctx,
emit_load_local_workgroup_id(struct ntd_context *ctx,
nir_intrinsic_instr *intr)
{
assert(intr->dest.is_ssa);
@ -3354,9 +3354,9 @@ emit_intrinsic(struct ntd_context *ctx, nir_intrinsic_instr *intr)
return emit_load_global_invocation_id(ctx, intr);
case nir_intrinsic_load_local_invocation_id:
return emit_load_local_invocation_id(ctx, intr);
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_work_group_id_zero_base:
return emit_load_local_work_group_id(ctx, intr);
case nir_intrinsic_load_workgroup_id:
case nir_intrinsic_load_workgroup_id_zero_base:
return emit_load_local_workgroup_id(ctx, intr);
case nir_intrinsic_load_ssbo:
return emit_load_ssbo(ctx, intr);
case nir_intrinsic_store_ssbo:
@ -3455,7 +3455,7 @@ emit_intrinsic(struct ntd_context *ctx, nir_intrinsic_instr *intr)
case nir_intrinsic_load_vulkan_descriptor:
return emit_load_vulkan_descriptor(ctx, intr);
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_num_workgroups:
case nir_intrinsic_load_workgroup_size:
default:
NIR_INSTR_UNSUPPORTED(&intr->instr);

4
src/panfrost/bifrost/bifrost_compile.c
View File

@ -1176,7 +1176,7 @@ bi_emit_intrinsic(bi_builder *b, nir_intrinsic_instr *instr)
case nir_intrinsic_load_viewport_scale:
case nir_intrinsic_load_viewport_offset:
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_num_workgroups:
case nir_intrinsic_load_workgroup_size:
bi_load_sysval_nir(b, instr, 3, 0);
break;
@ -1238,7 +1238,7 @@ bi_emit_intrinsic(bi_builder *b, nir_intrinsic_instr *instr)
bi_half(bi_register(55 + i / 2), i % 2));
break;
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_workgroup_id:
for (unsigned i = 0; i < 3; ++i)
bi_mov_i32_to(b, bi_word(dst, i), bi_register(57 + i));
break;

6
src/panfrost/midgard/midgard_compile.c
View File

@ -1515,7 +1515,7 @@ static unsigned
compute_builtin_arg(nir_intrinsic_op op)
{
switch (op) {
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_workgroup_id:
return REGISTER_LDST_GROUP_ID;
case nir_intrinsic_load_local_invocation_id:
return REGISTER_LDST_LOCAL_THREAD_ID;
@ -2023,13 +2023,13 @@ emit_intrinsic(compiler_context *ctx, nir_intrinsic_instr *instr)
case nir_intrinsic_load_viewport_scale:
case nir_intrinsic_load_viewport_offset:
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_num_workgroups:
case nir_intrinsic_load_sampler_lod_parameters_pan:
case nir_intrinsic_load_workgroup_size:
emit_sysval_read(ctx, &instr->instr, 3, 0);
break;
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_workgroup_id:
case nir_intrinsic_load_local_invocation_id:
case nir_intrinsic_load_global_invocation_id:
case nir_intrinsic_load_global_invocation_id_zero_base:

4
src/panfrost/util/pan_lower_64bit_intrin.c
View File

@ -44,8 +44,8 @@ nir_lower_64bit_intrin_instr(nir_builder *b, nir_instr *instr, void *data)
switch (intr->intrinsic) {
case nir_intrinsic_load_global_invocation_id:
case nir_intrinsic_load_global_invocation_id_zero_base:
case nir_intrinsic_load_work_group_id:
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_workgroup_id:
case nir_intrinsic_load_num_workgroups:
break;
default:

2
src/panfrost/util/pan_sysval.c
View File

@ -70,7 +70,7 @@ panfrost_nir_sysval_for_intrinsic(nir_intrinsic_instr *instr)
return PAN_SYSVAL_VIEWPORT_SCALE;
case nir_intrinsic_load_viewport_offset:
return PAN_SYSVAL_VIEWPORT_OFFSET;
case nir_intrinsic_load_num_work_groups:
case nir_intrinsic_load_num_workgroups:
return PAN_SYSVAL_NUM_WORK_GROUPS;
case nir_intrinsic_load_workgroup_size:
return PAN_SYSVAL_LOCAL_GROUP_SIZE;