Module: Mesa Branch: main Commit: 2fc2597fe533f4303fe6c1ab33f86a615730fb3b URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=2fc2597fe533f4303fe6c1ab33f86a615730fb3b
Author: Dave Airlie <[email protected]> Date: Wed Jul 5 10:22:06 2023 +1000 gallivm: make block_size use discrete values. Reviewed-by: Mike Blumenkrantz <[email protected]> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/23997> --- src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c | 9 ++++----- src/gallium/auxiliary/gallivm/lp_bld_tgsi.h | 2 +- src/gallium/drivers/llvmpipe/lp_state_cs.c | 7 +++---- 3 files changed, 8 insertions(+), 10 deletions(-) diff --git a/src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c b/src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c index 7eb32327801..408a305e000 100644 --- a/src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c +++ b/src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c @@ -1881,14 +1881,13 @@ static void emit_tex_size(struct lp_build_nir_context *bld_base, static LLVMValueRef get_local_invocation_index(struct lp_build_nir_soa_context *bld) { struct lp_build_nir_context *bld_base = &bld->bld_base; - struct gallivm_state *gallivm = bld_base->base.gallivm; LLVMValueRef tmp, tmp2; - tmp = lp_build_broadcast_scalar(&bld_base->uint_bld, LLVMBuildExtractElement(gallivm->builder, bld->system_values.block_size, lp_build_const_int32(gallivm, 1), "")); - tmp2 = lp_build_broadcast_scalar(&bld_base->uint_bld, LLVMBuildExtractElement(gallivm->builder, bld->system_values.block_size, lp_build_const_int32(gallivm, 0), "")); + + tmp = lp_build_broadcast_scalar(&bld_base->uint_bld, bld->system_values.block_size[1]); + tmp2 = lp_build_broadcast_scalar(&bld_base->uint_bld, bld->system_values.block_size[0]); tmp = lp_build_mul(&bld_base->uint_bld, tmp, tmp2); tmp = lp_build_mul(&bld_base->uint_bld, tmp, bld->system_values.thread_id[2]); - tmp2 = lp_build_broadcast_scalar(&bld_base->uint_bld, LLVMBuildExtractElement(gallivm->builder, bld->system_values.block_size, lp_build_const_int32(gallivm, 0), "")); tmp2 = lp_build_mul(&bld_base->uint_bld, tmp2, bld->system_values.thread_id[1]); tmp = lp_build_add(&bld_base->uint_bld, tmp, tmp2); tmp = lp_build_add(&bld_base->uint_bld, tmp, bld->system_values.thread_id[0]); @@ -1964,7 +1963,7 @@ static void emit_sysval_intrin(struct lp_build_nir_context *bld_base, break; 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), "")); + result[i] = lp_build_broadcast_scalar(&bld_base->uint_bld, bld->system_values.block_size[i]); break; case nir_intrinsic_load_work_dim: result[0] = lp_build_broadcast_scalar(&bld_base->uint_bld, bld->system_values.work_dim); diff --git a/src/gallium/auxiliary/gallivm/lp_bld_tgsi.h b/src/gallium/auxiliary/gallivm/lp_bld_tgsi.h index f48bb2f2f0c..0592a7f88dd 100644 --- a/src/gallium/auxiliary/gallivm/lp_bld_tgsi.h +++ b/src/gallium/auxiliary/gallivm/lp_bld_tgsi.h @@ -177,7 +177,7 @@ struct lp_bld_tgsi_system_values { LLVMValueRef grid_size[3]; LLVMValueRef front_facing; LLVMValueRef work_dim; - LLVMValueRef block_size; + LLVMValueRef block_size[3]; LLVMValueRef tess_coord; LLVMValueRef tess_outer; LLVMValueRef tess_inner; diff --git a/src/gallium/drivers/llvmpipe/lp_state_cs.c b/src/gallium/drivers/llvmpipe/lp_state_cs.c index 00cbeab1ab8..78e765e2aa4 100644 --- a/src/gallium/drivers/llvmpipe/lp_state_cs.c +++ b/src/gallium/drivers/llvmpipe/lp_state_cs.c @@ -692,10 +692,9 @@ generate_compute(struct llvmpipe_context *lp, LLVMValueRef subgroup_cmp = LLVMBuildICmp(gallivm->builder, LLVMIntEQ, num_subgroups, lp_build_const_int32(gallivm, 0), ""); system_values.num_subgroups = LLVMBuildSelect(builder, subgroup_cmp, lp_build_const_int32(gallivm, 1), num_subgroups, ""); - LLVMValueRef bsize[3] = { block_x_size_arg, block_y_size_arg, block_z_size_arg }; - system_values.block_size = LLVMGetUndef(LLVMVectorType(int32_type, 3)); - for (i = 0; i < 3; i++) - system_values.block_size = LLVMBuildInsertElement(builder, system_values.block_size, bsize[i], lp_build_const_int32(gallivm, i), ""); + system_values.block_size[0] = block_x_size_arg; + system_values.block_size[1] = block_y_size_arg; + system_values.block_size[2] = block_z_size_arg; LLVMValueRef last_x_loop = LLVMBuildICmp(gallivm->builder, LLVMIntEQ, x_size_arg, LLVMBuildSub(gallivm->builder, num_x_loop, lp_build_const_int32(gallivm, 1), ""), ""); LLVMValueRef use_partial_mask = LLVMBuildAnd(gallivm->builder, last_x_loop, has_partials, "");
