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

panvk: Automate sysval access from NIR shaders

Browse Source 
Emitting nir_load_push_constant() calls with the right base/range/offset
for sysvals is tedious and error prone. Provide syntactic sugar macros
to automate that.

Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com>
Reviewed-by: Chia-I Wu <olvaffe@gmail.com>
Reviewed-by: Mary Guillemard <mary.guillemard@collabora.com>
Reviewed-by: Lars-Ivar Hesselberg Simonsen <lars-ivar.simonsen@arm.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/32415>
This commit is contained in:
Boris Brezillon
2024-12-04 14:39:16 +01:00
committed by Marge Bot
parent cb20cb7b2f
commit 3d5ddaaffa
5 changed files with 83 additions and 71 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/panfrost/vulkan/jm/panvk_vX_cmd_dispatch.c
+2 -1
View File
@@ -90,7 +90,8 @@ panvk_per_arch(CmdDispatchBase)(VkCommandBuffer commandBuffer,
if (result != VK_SUCCESS)
return;
sysvals->desc.dyn_ssbos = cs_desc_state->dyn_ssbos;
sysvals->desc.sets[PANVK_DESC_TABLE_CS_DYN_SSBOS] =
cs_desc_state->dyn_ssbos;
}
for (uint32_t i = 0; i < MAX_SETS; i++) {
src/panfrost/vulkan/jm/panvk_vX_cmd_draw.c
+4 -2
View File
@@ -155,7 +155,8 @@ panvk_cmd_prepare_draw_sysvals(struct panvk_cmd_buffer *cmdbuf,
if (result != VK_SUCCESS)
return result;
sysvals->desc.vs_dyn_ssbos = vs_desc_state->dyn_ssbos;
sysvals->desc.sets[PANVK_DESC_TABLE_VS_DYN_SSBOS] =
vs_desc_state->dyn_ssbos;
gfx_state_set_dirty(cmdbuf, PUSH_UNIFORMS);
}
@@ -165,7 +166,8 @@ panvk_cmd_prepare_draw_sysvals(struct panvk_cmd_buffer *cmdbuf,
if (result != VK_SUCCESS)
return result;
sysvals->desc.fs_dyn_ssbos = fs_desc_state->dyn_ssbos;
sysvals->desc.sets[PANVK_DESC_TABLE_FS_DYN_SSBOS] =
fs_desc_state->dyn_ssbos;
gfx_state_set_dirty(cmdbuf, PUSH_UNIFORMS);
}
src/panfrost/vulkan/panvk_shader.h
+38 -5
View File
@@ -40,6 +40,17 @@ enum panvk_varying_buf_id {
PANVK_VARY_BUF_MAX,
};
#if PAN_ARCH <= 7
enum panvk_desc_table_id {
PANVK_DESC_TABLE_USER = 0,
PANVK_DESC_TABLE_CS_DYN_SSBOS = MAX_SETS,
PANVK_DESC_TABLE_COMPUTE_COUNT = PANVK_DESC_TABLE_CS_DYN_SSBOS + 1,
PANVK_DESC_TABLE_VS_DYN_SSBOS = MAX_SETS,
PANVK_DESC_TABLE_FS_DYN_SSBOS = MAX_SETS + 1,
PANVK_DESC_TABLE_GFX_COUNT = PANVK_DESC_TABLE_FS_DYN_SSBOS + 1,
};
#endif
struct panvk_graphics_sysvals {
struct {
struct {
@@ -65,9 +76,7 @@ struct panvk_graphics_sysvals {
int32_t layer_id;
struct {
uint64_t sets[MAX_SETS];
uint64_t vs_dyn_ssbos;
uint64_t fs_dyn_ssbos;
uint64_t sets[PANVK_DESC_TABLE_GFX_COUNT];
} desc;
#endif
};
@@ -85,14 +94,38 @@ struct panvk_compute_sysvals {
#if PAN_ARCH <= 7
struct {
uint64_t sets[MAX_SETS];
uint64_t dyn_ssbos;
uint64_t sets[PANVK_DESC_TABLE_COMPUTE_COUNT];
} desc;
#endif
};
#define SYSVALS_PUSH_CONST_BASE MAX_PUSH_CONSTANTS_SIZE
#define load_sysval(__b, __ptype, __bitsz, __name) \
nir_load_push_constant( \
__b, \
sizeof(((struct panvk_##__ptype##_sysvals *)NULL)->__name) / \
((__bitsz) / 8), \
__bitsz, \
nir_imm_int(__b, offsetof(struct panvk_##__ptype##_sysvals, __name)), \
.base = SYSVALS_PUSH_CONST_BASE, \
.range = sizeof(struct panvk_##__ptype##_sysvals))
#define load_sysval_entry(__b, __ptype, __bitsz, __name, __dyn_idx) \
nir_load_push_constant( \
__b, \
sizeof(((struct panvk_##__ptype##_sysvals *)NULL)->__name[0]) / \
((__bitsz) / 8), \
__bitsz, \
nir_iadd_imm( \
__b, \
nir_imul_imm( \
__b, __dyn_idx, \
sizeof(((struct panvk_##__ptype##_sysvals *)NULL)->__name[0])), \
offsetof(struct panvk_##__ptype##_sysvals, __name)), \
.base = SYSVALS_PUSH_CONST_BASE, \
.range = sizeof(struct panvk_##__ptype##_sysvals))
#if PAN_ARCH <= 7
enum panvk_bifrost_desc_table_type {
PANVK_BIFROST_DESC_TABLE_INVALID = -1,
src/panfrost/vulkan/panvk_vX_nir_lower_descriptors.c
+25 -25
View File
@@ -272,17 +272,20 @@ shader_ssbo_table(nir_builder *b, unsigned set, unsigned binding,
bind_layout->type == VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC);
bool is_dyn = bind_layout->type == VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC;
if (b->shader->info.stage == MESA_SHADER_COMPUTE)
return !is_dyn ? offsetof(struct panvk_compute_sysvals, desc.sets[set])
: offsetof(struct panvk_compute_sysvals, desc.dyn_ssbos);
else if (b->shader->info.stage == MESA_SHADER_VERTEX)
return !is_dyn
? offsetof(struct panvk_graphics_sysvals, desc.sets[set])
: offsetof(struct panvk_graphics_sysvals, desc.vs_dyn_ssbos);
else
return !is_dyn
? offsetof(struct panvk_graphics_sysvals, desc.sets[set])
: offsetof(struct panvk_graphics_sysvals, desc.fs_dyn_ssbos);
if (!is_dyn)
return PANVK_DESC_TABLE_USER + set;
switch (b->shader->info.stage) {
case MESA_SHADER_COMPUTE:
return PANVK_DESC_TABLE_CS_DYN_SSBOS;
case MESA_SHADER_VERTEX:
return PANVK_DESC_TABLE_VS_DYN_SSBOS;
case MESA_SHADER_FRAGMENT:
return PANVK_DESC_TABLE_FS_DYN_SSBOS;
default:
assert(!"Invalid stage");
return ~0;
}
}
#endif
@@ -331,9 +334,9 @@ build_res_index(nir_builder *b, uint32_t set, uint32_t binding,
case nir_address_format_64bit_bounded_global:
case nir_address_format_64bit_global_32bit_offset: {
unsigned base_addr_sysval_offs = shader_ssbo_table(b, set, binding, ctx);
unsigned desc_table = shader_ssbo_table(b, set, binding, ctx);
return nir_vec4(b, nir_imm_int(b, base_addr_sysval_offs),
return nir_vec4(b, nir_imm_int(b, desc_table),
nir_imm_int(b, desc_idx), array_index,
nir_imm_int(b, array_size - 1));
}
@@ -412,7 +415,7 @@ build_buffer_addr_for_res_index(nir_builder *b, nir_def *res_index,
case nir_address_format_64bit_bounded_global:
case nir_address_format_64bit_global_32bit_offset: {
nir_def *base_addr_sysval_offset = nir_channel(b, res_index, 0);
nir_def *desc_table_index = nir_channel(b, res_index, 0);
nir_def *first_desc_index = nir_channel(b, res_index, 1);
nir_def *array_index = nir_channel(b, res_index, 2);
nir_def *array_max = nir_channel(b, res_index, 3);
@@ -423,11 +426,11 @@ build_buffer_addr_for_res_index(nir_builder *b, nir_def *res_index,
nir_def *desc_offset = nir_imul_imm(
b, nir_iadd(b, array_index, first_desc_index), PANVK_DESCRIPTOR_SIZE);
nir_def *base_addr = nir_load_push_constant(
b, 1, 64, base_addr_sysval_offset, .base = SYSVALS_PUSH_CONST_BASE,
.range = b->shader->info.stage == MESA_SHADER_COMPUTE
? sizeof(struct panvk_compute_sysvals)
: sizeof(struct panvk_graphics_sysvals));
nir_def *base_addr =
b->shader->info.stage == MESA_SHADER_COMPUTE
? load_sysval_entry(b, compute, 64, desc.sets, desc_table_index)
: load_sysval_entry(b, graphics, 64, desc.sets, desc_table_index);
nir_def *desc_addr = nir_iadd(b, base_addr, nir_u2u64(b, desc_offset));
nir_def *desc =
nir_load_global(b, desc_addr, PANVK_DESCRIPTOR_SIZE, 4, 32);
@@ -557,13 +560,10 @@ load_resource_deref_desc(nir_builder *b, nir_deref_instr *deref,
set_offset = nir_iadd_imm(b, set_offset, desc_offset);
#if PAN_ARCH <= 7
unsigned set_base_addr_sysval_offs =
nir_def *set_base_addr =
b->shader->info.stage == MESA_SHADER_COMPUTE
? offsetof(struct panvk_compute_sysvals, desc.sets[set])
: offsetof(struct panvk_graphics_sysvals, desc.sets[set]);
nir_def *set_base_addr = nir_load_push_constant(
b, 1, 64, nir_imm_int(b, 0),
.base = SYSVALS_PUSH_CONST_BASE + set_base_addr_sysval_offs, .range = 8);
? load_sysval_entry(b, compute, 64, desc.sets, nir_imm_int(b, set))
: load_sysval_entry(b, graphics, 64, desc.sets, nir_imm_int(b, set));
unsigned desc_align = 1 << (ffs(PANVK_DESCRIPTOR_SIZE + desc_offset) - 1);
src/panfrost/vulkan/panvk_vX_shader.c
+14 -38
View File
@@ -56,17 +56,6 @@
#include "vk_shader.h"
#include "vk_util.h"
static nir_def *
load_sysval_from_push_const(nir_builder *b, unsigned offset, unsigned bit_size,
unsigned num_comps)
{
return nir_load_push_constant(
b, num_comps, bit_size, nir_imm_int(b, 0),
/* Push constants are placed first, and then come the sysvals. */
.base = offset + SYSVALS_PUSH_CONST_BASE,
.range = num_comps * bit_size / 8);
}
static bool
panvk_lower_sysvals(nir_builder *b, nir_instr *instr, void *data)
{
@@ -74,62 +63,49 @@ panvk_lower_sysvals(nir_builder *b, nir_instr *instr, void *data)
return false;
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
unsigned num_comps = intr->def.num_components;
unsigned bit_size = intr->def.bit_size;
nir_def *val = NULL;
b->cursor = nir_before_instr(instr);
#define SYSVAL(ptype, name) offsetof(struct panvk_##ptype##_sysvals, name)
switch (intr->intrinsic) {
case nir_intrinsic_load_base_workgroup_id:
val = load_sysval_from_push_const(b, SYSVAL(compute, base), bit_size,
num_comps);
val = load_sysval(b, compute, bit_size, base);
break;
case nir_intrinsic_load_num_workgroups:
val = load_sysval_from_push_const(b, SYSVAL(compute, num_work_groups),
bit_size, num_comps);
val = load_sysval(b, compute, bit_size, num_work_groups);
break;
case nir_intrinsic_load_workgroup_size:
val = load_sysval_from_push_const(b, SYSVAL(compute, local_group_size),
bit_size, num_comps);
val = load_sysval(b, compute, bit_size, local_group_size);
break;
case nir_intrinsic_load_viewport_scale:
val = load_sysval_from_push_const(b, SYSVAL(graphics, viewport.scale),
bit_size, num_comps);
val = load_sysval(b, graphics, bit_size, viewport.scale);
break;
case nir_intrinsic_load_viewport_offset:
val = load_sysval_from_push_const(b, SYSVAL(graphics, viewport.offset),
bit_size, num_comps);
val = load_sysval(b, graphics, bit_size, viewport.offset);
break;
case nir_intrinsic_load_first_vertex:
val = load_sysval_from_push_const(b, SYSVAL(graphics, vs.first_vertex),
bit_size, num_comps);
val = load_sysval(b, graphics, bit_size, vs.first_vertex);
break;
case nir_intrinsic_load_base_vertex:
val = load_sysval_from_push_const(b, SYSVAL(graphics, vs.base_vertex),
bit_size, num_comps);
val = load_sysval(b, graphics, bit_size, vs.base_vertex);
break;
case nir_intrinsic_load_base_instance:
val = load_sysval_from_push_const(b, SYSVAL(graphics, vs.base_instance),
bit_size, num_comps);
val = load_sysval(b, graphics, bit_size, vs.base_instance);
break;
case nir_intrinsic_load_blend_const_color_rgba:
val = load_sysval_from_push_const(b, SYSVAL(graphics, blend.constants),
bit_size, num_comps);
val = load_sysval(b, graphics, bit_size, blend.constants);
break;
case nir_intrinsic_load_noperspective_varyings_pan:
/* TODO: use a VS epilog specialized on constant noperspective_varyings
* with VK_EXT_graphics_pipeline_libraries and VK_EXT_shader_object */
assert(b->shader->info.stage == MESA_SHADER_VERTEX);
val = load_sysval_from_push_const(b, SYSVAL(graphics, vs.noperspective_varyings),
bit_size, num_comps);
val = load_sysval(b, graphics, bit_size, vs.noperspective_varyings);
break;
#if PAN_ARCH <= 7
case nir_intrinsic_load_layer_id:
assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
val = load_sysval_from_push_const(b, SYSVAL(graphics, layer_id), bit_size,
num_comps);
val = load_sysval(b, graphics, bit_size, layer_id);
break;
#endif
@@ -144,7 +120,8 @@ panvk_lower_sysvals(nir_builder *b, nir_instr *instr, void *data)
default:
return false;
}
#undef SYSVAL
assert(val->num_components == intr->def.num_components);
b->cursor = nir_after_instr(instr);
nir_def_rewrite_uses(&intr->def, val);
@@ -211,8 +188,7 @@ lower_gl_pos_layer_writes(nir_builder *b, nir_instr *instr, void *data)
nir_def *layer = nir_load_var(b, temp_layer_var);
nir_def *pos = nir_load_var(b, temp_pos_var);
nir_def *inf_pos = nir_imm_vec4(b, INFINITY, INFINITY, INFINITY, 1.0f);
nir_def *ref_layer = load_sysval_from_push_const(
b, offsetof(struct panvk_graphics_sysvals, layer_id), 32, 1);
nir_def *ref_layer = load_sysval(b, graphics, 32, layer_id);
nir_store_var(b, temp_pos_var,
nir_bcsel(b, nir_ieq(b, layer, ref_layer), pos, inf_pos),