diff --git a/src/panfrost/vulkan/jm/panvk_vX_cmd_dispatch.c b/src/panfrost/vulkan/jm/panvk_vX_cmd_dispatch.c index 31fac3c4c87..eb0cd547dec 100644 --- a/src/panfrost/vulkan/jm/panvk_vX_cmd_dispatch.c +++ b/src/panfrost/vulkan/jm/panvk_vX_cmd_dispatch.c @@ -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++) { diff --git a/src/panfrost/vulkan/jm/panvk_vX_cmd_draw.c b/src/panfrost/vulkan/jm/panvk_vX_cmd_draw.c index 76ae716c653..e3d2da86ad4 100644 --- a/src/panfrost/vulkan/jm/panvk_vX_cmd_draw.c +++ b/src/panfrost/vulkan/jm/panvk_vX_cmd_draw.c @@ -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); } diff --git a/src/panfrost/vulkan/panvk_shader.h b/src/panfrost/vulkan/panvk_shader.h index 9d0689fd7a3..40f851cc98e 100644 --- a/src/panfrost/vulkan/panvk_shader.h +++ b/src/panfrost/vulkan/panvk_shader.h @@ -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, diff --git a/src/panfrost/vulkan/panvk_vX_nir_lower_descriptors.c b/src/panfrost/vulkan/panvk_vX_nir_lower_descriptors.c index c9524bd472f..a97f87b32d4 100644 --- a/src/panfrost/vulkan/panvk_vX_nir_lower_descriptors.c +++ b/src/panfrost/vulkan/panvk_vX_nir_lower_descriptors.c @@ -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); diff --git a/src/panfrost/vulkan/panvk_vX_shader.c b/src/panfrost/vulkan/panvk_vX_shader.c index 15788b055fe..0c84e696232 100644 --- a/src/panfrost/vulkan/panvk_vX_shader.c +++ b/src/panfrost/vulkan/panvk_vX_shader.c @@ -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),