From: Caio Marcelo de Oliveira Filho Date: Thu, 27 May 2021 21:44:54 +0000 (-0700) Subject: nir: Rename nir_intrinsic_load_local_group_size to nir_intrinsic_load_workgroup_size X-Git-Tag: upstream/21.2.3~2366 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=a71a780598f598acea3efeb7fa7d05755dbcf0a8;p=platform%2Fupstream%2Fmesa.git nir: Rename nir_intrinsic_load_local_group_size to nir_intrinsic_load_workgroup_size Acked-by: Emma Anholt Acked-by: Alyssa Rosenzweig Reviewed-by: Jason Ekstrand Acked-by: Timur Kristóf Part-of: --- diff --git a/src/amd/llvm/ac_nir_to_llvm.c b/src/amd/llvm/ac_nir_to_llvm.c index be41ced..5fabf94 100644 --- a/src/amd/llvm/ac_nir_to_llvm.c +++ b/src/amd/llvm/ac_nir_to_llvm.c @@ -3417,7 +3417,7 @@ static void visit_intrinsic(struct ac_nir_context *ctx, nir_intrinsic_instr *ins result = ctx->abi->load_base_vertex(ctx->abi, instr->intrinsic == nir_intrinsic_load_base_vertex); break; - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: result = ctx->abi->load_local_group_size(ctx->abi); break; case nir_intrinsic_load_vertex_id: diff --git a/src/compiler/nir/nir.c b/src/compiler/nir/nir.c index 9aa91cd..2a78308 100644 --- a/src/compiler/nir/nir.c +++ b/src/compiler/nir/nir.c @@ -2020,7 +2020,7 @@ nir_intrinsic_from_system_value(gl_system_value val) case SYSTEM_VALUE_SUBGROUP_ID: return nir_intrinsic_load_subgroup_id; case SYSTEM_VALUE_WORKGROUP_SIZE: - return nir_intrinsic_load_local_group_size; + return nir_intrinsic_load_workgroup_size; case SYSTEM_VALUE_GLOBAL_INVOCATION_ID: return nir_intrinsic_load_global_invocation_id; case SYSTEM_VALUE_BASE_GLOBAL_INVOCATION_ID: @@ -2150,7 +2150,7 @@ nir_system_value_from_intrinsic(nir_intrinsic_op intrin) return SYSTEM_VALUE_NUM_SUBGROUPS; case nir_intrinsic_load_subgroup_id: return SYSTEM_VALUE_SUBGROUP_ID; - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: return SYSTEM_VALUE_WORKGROUP_SIZE; case nir_intrinsic_load_global_invocation_id: return SYSTEM_VALUE_GLOBAL_INVOCATION_ID; diff --git a/src/compiler/nir/nir_divergence_analysis.c b/src/compiler/nir/nir_divergence_analysis.c index 9777ee7..05d525f 100644 --- a/src/compiler/nir/nir_divergence_analysis.c +++ b/src/compiler/nir/nir_divergence_analysis.c @@ -102,7 +102,7 @@ visit_intrinsic(nir_shader *shader, nir_intrinsic_instr *instr) case nir_intrinsic_load_push_constant: case nir_intrinsic_load_work_dim: case nir_intrinsic_load_num_work_groups: - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: case nir_intrinsic_load_subgroup_id: case nir_intrinsic_load_num_subgroups: case nir_intrinsic_load_subgroup_size: diff --git a/src/compiler/nir/nir_gather_info.c b/src/compiler/nir/nir_gather_info.c index 5e8ad30..bdf84a3 100644 --- a/src/compiler/nir/nir_gather_info.c +++ b/src/compiler/nir/nir_gather_info.c @@ -625,7 +625,7 @@ gather_intrinsic_info(nir_intrinsic_instr *instr, nir_shader *shader, 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_local_group_size: + case nir_intrinsic_load_workgroup_size: case nir_intrinsic_load_work_dim: case nir_intrinsic_load_user_data_amd: case nir_intrinsic_load_view_index: diff --git a/src/compiler/nir/nir_intrinsics.py b/src/compiler/nir/nir_intrinsics.py index 18d24b9..89f8992 100644 --- a/src/compiler/nir/nir_intrinsics.py +++ b/src/compiler/nir/nir_intrinsics.py @@ -702,9 +702,9 @@ system_value("subgroup_le_mask", 0, bit_sizes=[32, 64]) system_value("subgroup_lt_mask", 0, bit_sizes=[32, 64]) system_value("num_subgroups", 1) system_value("subgroup_id", 1) -system_value("local_group_size", 3) +system_value("workgroup_size", 3) # note: the definition of global_invocation_id_zero_base is based on -# (work_group_id * local_group_size) + local_invocation_id. +# (work_group_id * workgroup_size) + local_invocation_id. # it is *not* based on work_group_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]) diff --git a/src/compiler/nir/nir_lower_system_values.c b/src/compiler/nir/nir_lower_system_values.c index 20e9603..5db6c2b 100644 --- a/src/compiler/nir/nir_lower_system_values.c +++ b/src/compiler/nir/nir_lower_system_values.c @@ -54,7 +54,7 @@ sanitize_32bit_sysval(nir_builder *b, nir_intrinsic_instr *intrin) static nir_ssa_def* build_global_group_size(nir_builder *b, unsigned bit_size) { - nir_ssa_def *group_size = nir_load_local_group_size(b); + nir_ssa_def *group_size = nir_load_workgroup_size(b); nir_ssa_def *num_work_groups = nir_load_num_work_groups(b, bit_size); return nir_imul(b, nir_u2u(b, group_size, bit_size), num_work_groups); @@ -116,7 +116,7 @@ lower_system_value_instr(nir_builder *b, nir_instr *instr, void *_state) case nir_intrinsic_load_local_invocation_id: case nir_intrinsic_load_local_invocation_index: - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: return sanitize_32bit_sysval(b, intrin); case nir_intrinsic_load_deref: { @@ -294,7 +294,7 @@ lower_compute_system_value_instr(nir_builder *b, * large so it can safely be omitted. */ nir_ssa_def *local_index = nir_load_local_invocation_index(b); - nir_ssa_def *local_size = nir_load_local_group_size(b); + nir_ssa_def *local_size = nir_load_workgroup_size(b); /* Because no hardware supports a local workgroup size greater than * about 1K, this calculation can be done in 32-bit and can save some @@ -324,7 +324,7 @@ lower_compute_system_value_instr(nir_builder *b, nir_ssa_def *size_x_imm; if (b->shader->info.cs.workgroup_size_variable) - size_x_imm = nir_channel(b, nir_load_local_group_size(b), 0); + size_x_imm = nir_channel(b, nir_load_workgroup_size(b), 0); else size_x_imm = nir_imm_int(b, size_x); @@ -424,7 +424,7 @@ lower_compute_system_value_instr(nir_builder *b, return NULL; } - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: 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 @@ -445,7 +445,7 @@ 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) || !b->shader->options->has_cs_global_id) { - nir_ssa_def *group_size = nir_load_local_group_size(b); + 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 *local_id = nir_load_local_invocation_id(b); diff --git a/src/freedreno/ir3/ir3_compiler_nir.c b/src/freedreno/ir3/ir3_compiler_nir.c index 195b38f..1ff3aa3 100644 --- a/src/freedreno/ir3/ir3_compiler_nir.c +++ b/src/freedreno/ir3/ir3_compiler_nir.c @@ -2042,7 +2042,7 @@ emit_intrinsic(struct ir3_context *ctx, nir_intrinsic_instr *intr) dst[i] = create_driver_param(ctx, IR3_DP_NUM_WORK_GROUPS_X + i); } break; - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: for (int i = 0; i < dest_components; i++) { dst[i] = create_driver_param(ctx, IR3_DP_LOCAL_GROUP_SIZE_X + i); } diff --git a/src/freedreno/ir3/ir3_nir.c b/src/freedreno/ir3/ir3_nir.c index c8d0a0b..0353bb3 100644 --- a/src/freedreno/ir3/ir3_nir.c +++ b/src/freedreno/ir3/ir3_nir.c @@ -680,7 +680,7 @@ ir3_nir_scan_driver_consts(nir_shader *shader, layout->num_driver_params = MAX2(layout->num_driver_params, IR3_DP_NUM_WORK_GROUPS_Z + 1); break; - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: layout->num_driver_params = MAX2(layout->num_driver_params, IR3_DP_LOCAL_GROUP_SIZE_Z + 1); break; diff --git a/src/gallium/auxiliary/gallivm/lp_bld_nir.c b/src/gallium/auxiliary/gallivm/lp_bld_nir.c index 5d2c27e..bab9975 100644 --- a/src/gallium/auxiliary/gallivm/lp_bld_nir.c +++ b/src/gallium/auxiliary/gallivm/lp_bld_nir.c @@ -1696,7 +1696,7 @@ static void visit_intrinsic(struct lp_build_nir_context *bld_base, case nir_intrinsic_load_invocation_id: case nir_intrinsic_load_front_face: case nir_intrinsic_load_draw_id: - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: case nir_intrinsic_load_work_dim: case nir_intrinsic_load_tess_coord: case nir_intrinsic_load_tess_level_outer: diff --git a/src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c b/src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c index 05e5208..45f9957 100644 --- a/src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c +++ b/src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c @@ -1552,7 +1552,7 @@ static void emit_sysval_intrin(struct lp_build_nir_context *bld_base, break; default: break; - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: for (unsigned i = 0; i < 3; i++) result[i] = lp_build_broadcast_scalar(&bld_base->uint_bld, LLVMBuildExtractElement(gallivm->builder, bld->system_values.block_size, lp_build_const_int32(gallivm, i), "")); break; diff --git a/src/gallium/auxiliary/nir/nir_to_tgsi.c b/src/gallium/auxiliary/nir/nir_to_tgsi.c index 3c73d34..69fd597 100644 --- a/src/gallium/auxiliary/nir/nir_to_tgsi.c +++ b/src/gallium/auxiliary/nir/nir_to_tgsi.c @@ -1652,7 +1652,7 @@ ntt_emit_intrinsic(struct ntt_compile *c, nir_intrinsic_instr *instr) 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_local_group_size: + case nir_intrinsic_load_workgroup_size: case nir_intrinsic_load_subgroup_size: case nir_intrinsic_load_subgroup_invocation: case nir_intrinsic_load_subgroup_eq_mask: diff --git a/src/gallium/auxiliary/nir/nir_to_tgsi_info.c b/src/gallium/auxiliary/nir/nir_to_tgsi_info.c index e74e90b..8b1a9e7 100644 --- a/src/gallium/auxiliary/nir/nir_to_tgsi_info.c +++ b/src/gallium/auxiliary/nir/nir_to_tgsi_info.c @@ -223,7 +223,7 @@ static void scan_instruction(const struct nir_shader *nir, case nir_intrinsic_load_num_work_groups: info->uses_grid_size = true; break; - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: /* The block size is translated to IMM with a fixed block size. */ if (info->properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] == 0) info->uses_block_size = true; diff --git a/src/gallium/auxiliary/nir/tgsi_to_nir.c b/src/gallium/auxiliary/nir/tgsi_to_nir.c index 0a7de87..339d1cc 100644 --- a/src/gallium/auxiliary/nir/tgsi_to_nir.c +++ b/src/gallium/auxiliary/nir/tgsi_to_nir.c @@ -623,7 +623,7 @@ ttn_src_for_file_and_index(struct ttn_compile *c, unsigned file, unsigned index, load = nir_load_work_group_id(b, 32); break; case TGSI_SEMANTIC_BLOCK_SIZE: - load = nir_load_local_group_size(b); + load = nir_load_workgroup_size(b); break; case TGSI_SEMANTIC_CS_USER_DATA_AMD: load = nir_load_user_data_amd(b); diff --git a/src/gallium/drivers/iris/iris_program.c b/src/gallium/drivers/iris/iris_program.c index 92915a9..9028085 100644 --- a/src/gallium/drivers/iris/iris_program.c +++ b/src/gallium/drivers/iris/iris_program.c @@ -527,7 +527,7 @@ iris_setup_uniforms(const struct brw_compiler *compiler, nir_intrinsic_base(intrin) * 16)); break; } - case nir_intrinsic_load_local_group_size: { + case nir_intrinsic_load_workgroup_size: { assert(nir->info.cs.workgroup_size_variable); if (variable_group_size_idx == -1) { variable_group_size_idx = num_system_values; diff --git a/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp b/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp index ad7d104..62c544b 100644 --- a/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp +++ b/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp @@ -1566,7 +1566,7 @@ Converter::convert(nir_intrinsic_op intr) return SV_INSTANCE_ID; case nir_intrinsic_load_invocation_id: return SV_INVOCATION_ID; - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: return SV_NTID; case nir_intrinsic_load_local_invocation_id: return SV_TID; @@ -1843,7 +1843,7 @@ Converter::visit(nir_intrinsic_instr *insn) case nir_intrinsic_load_helper_invocation: case nir_intrinsic_load_instance_id: case nir_intrinsic_load_invocation_id: - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: case nir_intrinsic_load_local_invocation_id: case nir_intrinsic_load_num_work_groups: case nir_intrinsic_load_patch_vertices_in: diff --git a/src/gallium/drivers/radeonsi/si_shaderlib_nir.c b/src/gallium/drivers/radeonsi/si_shaderlib_nir.c index ab43b8e..8b9fd35 100644 --- a/src/gallium/drivers/radeonsi/si_shaderlib_nir.c +++ b/src/gallium/drivers/radeonsi/si_shaderlib_nir.c @@ -43,7 +43,7 @@ static nir_ssa_def *get_global_ids(nir_builder *b, unsigned 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_size = nir_channels(b, nir_load_local_group_size(b), 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); } diff --git a/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c b/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c index a248379..ab7eaa2 100644 --- a/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c +++ b/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c @@ -2785,7 +2785,7 @@ emit_intrinsic(struct ntv_context *ctx, nir_intrinsic_instr *intr) emit_load_uint_input(ctx, intr, &ctx->local_invocation_index_var, "gl_LocalInvocationIndex", SpvBuiltInLocalInvocationIndex); break; - case nir_intrinsic_load_local_group_size: { + case nir_intrinsic_load_workgroup_size: { assert(ctx->local_group_size_var); store_dest(ctx, &intr->dest, ctx->local_group_size_var, nir_type_uint); break; diff --git a/src/intel/compiler/brw_fs_nir.cpp b/src/intel/compiler/brw_fs_nir.cpp index eac4429..d50b7ab 100644 --- a/src/intel/compiler/brw_fs_nir.cpp +++ b/src/intel/compiler/brw_fs_nir.cpp @@ -3814,7 +3814,7 @@ fs_visitor::nir_emit_cs_intrinsic(const fs_builder &bld, break; } - case nir_intrinsic_load_local_group_size: { + case nir_intrinsic_load_workgroup_size: { assert(compiler->lower_variable_group_size); assert(nir->info.cs.workgroup_size_variable); for (unsigned i = 0; i < 3; i++) { diff --git a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c index 5dc6dd7..66999f4 100644 --- a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c +++ b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c @@ -53,7 +53,7 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state, nir_ssa_def *sysval; switch (intrinsic->intrinsic) { - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: case nir_intrinsic_load_work_group_id: case nir_intrinsic_load_num_work_groups: /* Convert this to 32-bit if it's not */ @@ -82,7 +82,7 @@ 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.workgroup_size_variable) { - nir_ssa_def *size_xyz = nir_load_local_group_size(b); + nir_ssa_def *size_xyz = nir_load_workgroup_size(b); size_x = nir_channel(b, size_xyz, 0); size_y = nir_channel(b, size_xyz, 1); } else { @@ -214,7 +214,7 @@ 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.workgroup_size_variable) { - nir_ssa_def *size_xyz = nir_load_local_group_size(b); + nir_ssa_def *size_xyz = nir_load_workgroup_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); diff --git a/src/microsoft/clc/clc_nir.c b/src/microsoft/clc/clc_nir.c index eaac4e6..3d8fb7c 100644 --- a/src/microsoft/clc/clc_nir.c +++ b/src/microsoft/clc/clc_nir.c @@ -146,7 +146,7 @@ clc_nir_lower_system_values(nir_shader *nir, nir_variable *var) case nir_intrinsic_load_work_dim: progress |= lower_load_work_dim(&b, intr, var); break; - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: lower_load_local_group_size(&b, intr); break; case nir_intrinsic_load_num_work_groups: diff --git a/src/microsoft/compiler/nir_to_dxil.c b/src/microsoft/compiler/nir_to_dxil.c index 3ea4183..6c16268 100644 --- a/src/microsoft/compiler/nir_to_dxil.c +++ b/src/microsoft/compiler/nir_to_dxil.c @@ -3456,7 +3456,7 @@ emit_intrinsic(struct ntd_context *ctx, nir_intrinsic_instr *intr) return emit_load_vulkan_descriptor(ctx, intr); case nir_intrinsic_load_num_work_groups: - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: default: NIR_INSTR_UNSUPPORTED(&intr->instr); assert("Unimplemented intrinsic instruction"); diff --git a/src/panfrost/bifrost/bifrost_compile.c b/src/panfrost/bifrost/bifrost_compile.c index d73987d..9189618 100644 --- a/src/panfrost/bifrost/bifrost_compile.c +++ b/src/panfrost/bifrost/bifrost_compile.c @@ -1177,7 +1177,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_local_group_size: + case nir_intrinsic_load_workgroup_size: bi_load_sysval_nir(b, instr, 3, 0); break; diff --git a/src/panfrost/midgard/midgard_compile.c b/src/panfrost/midgard/midgard_compile.c index c04d591..d0ae5a7 100644 --- a/src/panfrost/midgard/midgard_compile.c +++ b/src/panfrost/midgard/midgard_compile.c @@ -2025,7 +2025,7 @@ emit_intrinsic(compiler_context *ctx, nir_intrinsic_instr *instr) case nir_intrinsic_load_viewport_offset: case nir_intrinsic_load_num_work_groups: case nir_intrinsic_load_sampler_lod_parameters_pan: - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: emit_sysval_read(ctx, &instr->instr, 3, 0); break; diff --git a/src/panfrost/util/pan_sysval.c b/src/panfrost/util/pan_sysval.c index fd261a8..0715058 100644 --- a/src/panfrost/util/pan_sysval.c +++ b/src/panfrost/util/pan_sysval.c @@ -72,7 +72,7 @@ panfrost_nir_sysval_for_intrinsic(nir_intrinsic_instr *instr) return PAN_SYSVAL_VIEWPORT_OFFSET; case nir_intrinsic_load_num_work_groups: return PAN_SYSVAL_NUM_WORK_GROUPS; - case nir_intrinsic_load_local_group_size: + case nir_intrinsic_load_workgroup_size: return PAN_SYSVAL_LOCAL_GROUP_SIZE; case nir_intrinsic_load_work_dim: return PAN_SYSVAL_WORK_DIM;