fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-04 11:58:10 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 2

nir/opt_preamble: make load_workgroup_size handling optional

Browse source 
not all drivers support it being in the preamble, e.g. asahi.

Signed-off-by: Karol Herbst <git@karolherbst.de>
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26614>
This commit is contained in:
Karol Herbst 2023-08-29 20:54:14 +02:00 • committed by Alyssa Rosenzweig
parent 3ddb07c36b
commit 6979a1aa07
4 changed files with 9 additions and 1 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

2
src/asahi/compiler/agx_nir_opt_preamble.c
View file

@ -116,6 +116,8 @@ avoid_instr(const nir_instr *instr, const void *data)
static const nir_opt_preamble_options preamble_options = {
.drawid_uniform = true,
.subgroup_size_uniform = true,
/* not supported in hardware */
.load_workgroup_size_allowed = false,
.def_size = def_size,
.instr_cost_cb = instr_cost,
.rewrite_cost_cb = rewrite_cost,

3
src/compiler/nir/nir.h
View file

@ -6455,6 +6455,9 @@ typedef struct {
/* True if the subgroup size is uniform. */
bool subgroup_size_uniform;
/* True if load_workgroup_size is supported in the preamble. */
bool load_workgroup_size_allowed;
/* size/align for load/store_preamble. */
void (*def_size)(nir_def *def, unsigned *size, unsigned *align);

4
src/compiler/nir/nir_opt_preamble.c
View file

@ -150,7 +150,6 @@ can_move_intrinsic(nir_intrinsic_instr *instr, opt_preamble_ctx *ctx)
case nir_intrinsic_load_push_constant:
case nir_intrinsic_load_work_dim:
case nir_intrinsic_load_num_workgroups:
case nir_intrinsic_load_workgroup_size:
case nir_intrinsic_load_ray_launch_size:
case nir_intrinsic_load_ray_launch_size_addr_amd:
case nir_intrinsic_load_sbt_base_amd:
@ -200,6 +199,9 @@ can_move_intrinsic(nir_intrinsic_instr *instr, opt_preamble_ctx *ctx)
case nir_intrinsic_load_num_subgroups:
return ctx->options->subgroup_size_uniform;
case nir_intrinsic_load_workgroup_size:
return ctx->options->load_workgroup_size_allowed;
/* Intrinsics which can be moved if the sources can */
case nir_intrinsic_load_ubo:
case nir_intrinsic_load_ubo_vec4:

1
src/freedreno/ir3/ir3_nir_opt_preamble.c
View file

@ -319,6 +319,7 @@ ir3_nir_opt_preamble(nir_shader *nir, struct ir3_shader_variant *v)
nir_opt_preamble_options options = {
.drawid_uniform = true,
.subgroup_size_uniform = true,
.load_workgroup_size_allowed = true,
.def_size = def_size,
.preamble_storage_size = max_size,
.instr_cost_cb = instr_cost,