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

gallivm: make block_size use discrete values.

Browse Source 
Reviewed-by: Mike Blumenkrantz <michael.blumenkrantz@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/23997>
This commit is contained in:
Dave Airlie
2023-07-05 10:22:06 +10:00
parent c0123d792c
commit 2fc2597fe5
3 changed files with 8 additions and 10 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c
+4 -5
View File
@@ -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);
src/gallium/auxiliary/gallivm/lp_bld_tgsi.h
+1 -1
View File
@@ -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;
src/gallium/drivers/llvmpipe/lp_state_cs.c
+3 -4
View File
@@ -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, "");