diff --git a/src/gallium/drivers/panfrost/pan_csf.c b/src/gallium/drivers/panfrost/pan_csf.c index 0e470e4463c..d5cd8aa3673 100644 --- a/src/gallium/drivers/panfrost/pan_csf.c +++ b/src/gallium/drivers/panfrost/pan_csf.c @@ -881,10 +881,10 @@ GENX(csf_launch_grid)(struct panfrost_batch *batch, csf_emit_shader_regs(batch, PIPE_SHADER_COMPUTE, batch->rsd[PIPE_SHADER_COMPUTE]); - cs_move64_to(b, cs_reg64(b, 24), batch->tls.gpu); + cs_move64_to(b, cs_reg64(b, MALI_COMPUTE_SR_TSD_0), batch->tls.gpu); /* Global attribute offset */ - cs_move32_to(b, cs_reg32(b, 32), 0); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_GLOBAL_ATTRIBUTE_OFFSET), 0); /* Compute workgroup size */ struct mali_compute_size_workgroup_packed wg_size; @@ -903,11 +903,11 @@ GENX(csf_launch_grid)(struct panfrost_batch *batch, (info->variable_shared_mem == 0); } - cs_move32_to(b, cs_reg32(b, 33), wg_size.opaque[0]); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_WG_SIZE), wg_size.opaque[0]); - /* Offset */ - for (unsigned i = 0; i < 3; ++i) - cs_move32_to(b, cs_reg32(b, 34 + i), 0); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_OFFSET_X), 0); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_OFFSET_Y), 0); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_OFFSET_Z), 0); unsigned threads_per_wg = info->block[0] * info->block[1] * info->block[2]; unsigned max_thread_cnt = panfrost_compute_max_thread_count( @@ -920,7 +920,7 @@ GENX(csf_launch_grid)(struct panfrost_batch *batch, b, address, pan_resource(info->indirect)->image.data.base + info->indirect_offset); - struct cs_index grid_xyz = cs_reg_tuple(b, 37, 3); + struct cs_index grid_xyz = cs_reg_tuple(b, MALI_COMPUTE_SR_JOB_SIZE_X, 3); cs_load_to(b, grid_xyz, address, BITFIELD_MASK(3), 0); /* Wait for the load */ @@ -942,8 +942,9 @@ GENX(csf_launch_grid)(struct panfrost_batch *batch, false, cs_shader_res_sel(0, 0, 0, 0)); } else { /* Set size in workgroups per dimension immediately */ - for (unsigned i = 0; i < 3; ++i) - cs_move32_to(b, cs_reg32(b, 37 + i), info->grid[i]); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_X), info->grid[0]); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_Y), info->grid[1]); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_Z), info->grid[2]); /* Pick the task_axis and task_increment to maximize thread utilization. */ unsigned task_axis = MALI_TASK_AXIS_X; @@ -984,10 +985,11 @@ GENX(csf_launch_xfb)(struct panfrost_batch *batch, { struct cs_builder *b = batch->csf.cs.builder; - cs_move64_to(b, cs_reg64(b, 24), batch->tls.gpu); + cs_move64_to(b, cs_reg64(b, MALI_COMPUTE_SR_TSD_0), batch->tls.gpu); /* TODO: Indexing. Also, attribute_offset is a legacy feature.. */ - cs_move32_to(b, cs_reg32(b, 32), batch->ctx->offset_start); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_GLOBAL_ATTRIBUTE_OFFSET), + batch->ctx->offset_start); /* Compute workgroup size */ struct mali_compute_size_workgroup_packed wg_size; @@ -1001,15 +1003,16 @@ GENX(csf_launch_xfb)(struct panfrost_batch *batch, */ cfg.allow_merging_workgroups = true; } - cs_move32_to(b, cs_reg32(b, 33), wg_size.opaque[0]); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_WG_SIZE), wg_size.opaque[0]); - /* Offset */ - for (unsigned i = 0; i < 3; ++i) - cs_move32_to(b, cs_reg32(b, 34 + i), 0); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_OFFSET_X), 0); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_OFFSET_Y), 0); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_OFFSET_Z), 0); - cs_move32_to(b, cs_reg32(b, 37), count); - cs_move32_to(b, cs_reg32(b, 38), info->instance_count); - cs_move32_to(b, cs_reg32(b, 39), 1); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_X), count); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_Y), + info->instance_count); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_Z), 1); csf_emit_shader_regs(batch, PIPE_SHADER_VERTEX, batch->rsd[PIPE_SHADER_VERTEX]); diff --git a/src/gallium/drivers/panfrost/pan_precomp.c b/src/gallium/drivers/panfrost/pan_precomp.c index dd0c01f668c..7de146f6be1 100644 --- a/src/gallium/drivers/panfrost/pan_precomp.c +++ b/src/gallium/drivers/panfrost/pan_precomp.c @@ -320,18 +320,18 @@ GENX(panfrost_launch_precomp)(struct panfrost_batch *batch, struct cs_builder *b = batch->csf.cs.builder; /* No resource table */ - cs_move64_to(b, cs_reg64(b, 0), 0); + cs_move64_to(b, cs_reg64(b, MALI_COMPUTE_SR_SRT_0), 0); uint64_t fau_count = DIV_ROUND_UP(BIFROST_PRECOMPILED_KERNEL_SYSVALS_SIZE + data_size, 8); uint64_t fau_ptr = push_uniforms.gpu | (fau_count << 56); - cs_move64_to(b, cs_reg64(b, 8), fau_ptr); + cs_move64_to(b, cs_reg64(b, MALI_COMPUTE_SR_FAU_0), fau_ptr); - cs_move64_to(b, cs_reg64(b, 16), shader->state_ptr); - cs_move64_to(b, cs_reg64(b, 24), tsd); + cs_move64_to(b, cs_reg64(b, MALI_COMPUTE_SR_SPD_0), shader->state_ptr); + cs_move64_to(b, cs_reg64(b, MALI_COMPUTE_SR_TSD_0), tsd); /* Global attribute offset */ - cs_move32_to(b, cs_reg32(b, 32), 0); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_GLOBAL_ATTRIBUTE_OFFSET), 0); /* Compute workgroup size */ struct mali_compute_size_workgroup_packed wg_size; @@ -341,17 +341,17 @@ GENX(panfrost_launch_precomp)(struct panfrost_batch *batch, cfg.workgroup_size_z = shader->local_size.z; cfg.allow_merging_workgroups = false; } - cs_move32_to(b, cs_reg32(b, 33), wg_size.opaque[0]); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_WG_SIZE), wg_size.opaque[0]); /* Job offset */ - cs_move32_to(b, cs_reg32(b, 34), 0); - cs_move32_to(b, cs_reg32(b, 35), 0); - cs_move32_to(b, cs_reg32(b, 36), 0); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_OFFSET_X), 0); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_OFFSET_Y), 0); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_OFFSET_Z), 0); /* Job size */ - cs_move32_to(b, cs_reg32(b, 37), grid.count[0]); - cs_move32_to(b, cs_reg32(b, 38), grid.count[1]); - cs_move32_to(b, cs_reg32(b, 39), grid.count[2]); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_X), grid.count[0]); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_Y), grid.count[1]); + cs_move32_to(b, cs_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_Z), grid.count[2]); unsigned threads_per_wg = shader->local_size.x * shader->local_size.y * shader->local_size.z; diff --git a/src/panfrost/lib/genxml/v10.xml b/src/panfrost/lib/genxml/v10.xml index bf395ecf001..e97d6d98625 100644 --- a/src/panfrost/lib/genxml/v10.xml +++ b/src/panfrost/lib/genxml/v10.xml @@ -840,6 +840,33 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/src/panfrost/vulkan/csf/panvk_vX_cmd_dispatch.c b/src/panfrost/vulkan/csf/panvk_vX_cmd_dispatch.c index b8511557abc..3f67b9ad423 100644 --- a/src/panfrost/vulkan/csf/panvk_vX_cmd_dispatch.c +++ b/src/panfrost/vulkan/csf/panvk_vX_cmd_dispatch.c @@ -236,22 +236,24 @@ cmd_dispatch(struct panvk_cmd_buffer *cmdbuf, struct panvk_dispatch_info *info) cs_update_compute_ctx(b) { if (compute_state_dirty(cmdbuf, CS) || compute_state_dirty(cmdbuf, DESC_STATE)) - cs_move64_to(b, cs_sr_reg64(b, 0), cs_desc_state->res_table); + cs_move64_to(b, cs_sr_reg64(b, MALI_COMPUTE_SR_SRT_0), + cs_desc_state->res_table); if (compute_state_dirty(cmdbuf, PUSH_UNIFORMS)) { uint64_t fau_ptr = cmdbuf->state.compute.push_uniforms | ((uint64_t)shader->fau.total_count << 56); - cs_move64_to(b, cs_sr_reg64(b, 8), fau_ptr); + cs_move64_to(b, cs_sr_reg64(b, MALI_COMPUTE_SR_FAU_0), fau_ptr); } if (compute_state_dirty(cmdbuf, CS)) - cs_move64_to(b, cs_sr_reg64(b, 16), + cs_move64_to(b, cs_sr_reg64(b, MALI_COMPUTE_SR_SPD_0), panvk_priv_mem_dev_addr(shader->spd)); - cs_move64_to(b, cs_sr_reg64(b, 24), tsd); + cs_move64_to(b, cs_sr_reg64(b, MALI_COMPUTE_SR_TSD_0), tsd); /* Global attribute offset */ - cs_move32_to(b, cs_sr_reg32(b, 32), 0); + cs_move32_to(b, cs_sr_reg32(b, MALI_COMPUTE_SR_GLOBAL_ATTRIBUTE_OFFSET), + 0); struct mali_compute_size_workgroup_packed wg_size; pan_pack(&wg_size, COMPUTE_SIZE_WORKGROUP, cfg) { @@ -260,47 +262,54 @@ cmd_dispatch(struct panvk_cmd_buffer *cmdbuf, struct panvk_dispatch_info *info) cfg.workgroup_size_z = shader->local_size.z; cfg.allow_merging_workgroups = false; } - cs_move32_to(b, cs_sr_reg32(b, 33), wg_size.opaque[0]); - cs_move32_to(b, cs_sr_reg32(b, 34), + cs_move32_to(b, cs_sr_reg32(b, MALI_COMPUTE_SR_WG_SIZE), + wg_size.opaque[0]); + cs_move32_to(b, cs_sr_reg32(b, MALI_COMPUTE_SR_JOB_OFFSET_X), info->wg_base.x * shader->local_size.x); - cs_move32_to(b, cs_sr_reg32(b, 35), + cs_move32_to(b, cs_sr_reg32(b, MALI_COMPUTE_SR_JOB_OFFSET_Y), info->wg_base.y * shader->local_size.y); - cs_move32_to(b, cs_sr_reg32(b, 36), + cs_move32_to(b, cs_sr_reg32(b, MALI_COMPUTE_SR_JOB_OFFSET_Z), info->wg_base.z * shader->local_size.z); if (indirect) { /* Load parameters from indirect buffer and update workgroup count * registers and sysvals */ cs_move64_to(b, cs_scratch_reg64(b, 0), info->indirect.buffer_dev_addr); - cs_load_to(b, cs_sr_reg_tuple(b, 37, 3), cs_scratch_reg64(b, 0), - BITFIELD_MASK(3), 0); + cs_load_to(b, cs_sr_reg_tuple(b, MALI_COMPUTE_SR_JOB_SIZE_X, 3), + cs_scratch_reg64(b, 0), BITFIELD_MASK(3), 0); cs_move64_to(b, cs_scratch_reg64(b, 0), cmdbuf->state.compute.push_uniforms); cs_wait_slot(b, SB_ID(LS), false); if (shader_uses_sysval(shader, compute, num_work_groups.x)) { - cs_store32(b, cs_sr_reg32(b, 37), cs_scratch_reg64(b, 0), + cs_store32(b, cs_sr_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_X), + cs_scratch_reg64(b, 0), shader_remapped_sysval_offset( shader, sysval_offset(compute, num_work_groups.x))); } if (shader_uses_sysval(shader, compute, num_work_groups.y)) { - cs_store32(b, cs_sr_reg32(b, 38), cs_scratch_reg64(b, 0), + cs_store32(b, cs_sr_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_Y), + cs_scratch_reg64(b, 0), shader_remapped_sysval_offset( shader, sysval_offset(compute, num_work_groups.y))); } if (shader_uses_sysval(shader, compute, num_work_groups.z)) { - cs_store32(b, cs_sr_reg32(b, 39), cs_scratch_reg64(b, 0), + cs_store32(b, cs_sr_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_Z), + cs_scratch_reg64(b, 0), shader_remapped_sysval_offset( shader, sysval_offset(compute, num_work_groups.z))); } cs_wait_slot(b, SB_ID(LS), false); } else { - cs_move32_to(b, cs_sr_reg32(b, 37), info->direct.wg_count.x); - cs_move32_to(b, cs_sr_reg32(b, 38), info->direct.wg_count.y); - cs_move32_to(b, cs_sr_reg32(b, 39), info->direct.wg_count.z); + cs_move32_to(b, cs_sr_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_X), + info->direct.wg_count.x); + cs_move32_to(b, cs_sr_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_Y), + info->direct.wg_count.y); + cs_move32_to(b, cs_sr_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_Z), + info->direct.wg_count.z); } } diff --git a/src/panfrost/vulkan/csf/panvk_vX_cmd_precomp.c b/src/panfrost/vulkan/csf/panvk_vX_cmd_precomp.c index dfed6281d1f..2e88e01dd36 100644 --- a/src/panfrost/vulkan/csf/panvk_vX_cmd_precomp.c +++ b/src/panfrost/vulkan/csf/panvk_vX_cmd_precomp.c @@ -103,19 +103,21 @@ panvk_per_arch(dispatch_precomp)(struct panvk_precomp_ctx *ctx, cs_update_compute_ctx(b) { /* No resource table */ - cs_move64_to(b, cs_sr_reg64(b, 0), 0); + cs_move64_to(b, cs_sr_reg64(b, MALI_COMPUTE_SR_SRT_0), 0); uint64_t fau_count = DIV_ROUND_UP(BIFROST_PRECOMPILED_KERNEL_SYSVALS_SIZE + data_size, 8); uint64_t fau_ptr = push_uniforms.gpu | (fau_count << 56); - cs_move64_to(b, cs_sr_reg64(b, 8), fau_ptr); + cs_move64_to(b, cs_sr_reg64(b, MALI_COMPUTE_SR_FAU_0), fau_ptr); - cs_move64_to(b, cs_sr_reg64(b, 16), panvk_priv_mem_dev_addr(shader->spd)); + cs_move64_to(b, cs_sr_reg64(b, MALI_COMPUTE_SR_SPD_0), + panvk_priv_mem_dev_addr(shader->spd)); - cs_move64_to(b, cs_sr_reg64(b, 24), tsd); + cs_move64_to(b, cs_sr_reg64(b, MALI_COMPUTE_SR_TSD_0), tsd); /* Global attribute offset */ - cs_move32_to(b, cs_sr_reg32(b, 32), 0); + cs_move32_to(b, cs_sr_reg32(b, MALI_COMPUTE_SR_GLOBAL_ATTRIBUTE_OFFSET), + 0); struct mali_compute_size_workgroup_packed wg_size; pan_pack(&wg_size, COMPUTE_SIZE_WORKGROUP, cfg) { @@ -124,17 +126,21 @@ panvk_per_arch(dispatch_precomp)(struct panvk_precomp_ctx *ctx, cfg.workgroup_size_z = shader->local_size.z; cfg.allow_merging_workgroups = false; } - cs_move32_to(b, cs_sr_reg32(b, 33), wg_size.opaque[0]); + cs_move32_to(b, cs_sr_reg32(b, MALI_COMPUTE_SR_WG_SIZE), + wg_size.opaque[0]); /* Job offset */ - cs_move32_to(b, cs_sr_reg32(b, 34), 0); - cs_move32_to(b, cs_sr_reg32(b, 35), 0); - cs_move32_to(b, cs_sr_reg32(b, 36), 0); + cs_move32_to(b, cs_sr_reg32(b, MALI_COMPUTE_SR_JOB_OFFSET_X), 0); + cs_move32_to(b, cs_sr_reg32(b, MALI_COMPUTE_SR_JOB_OFFSET_Y), 0); + cs_move32_to(b, cs_sr_reg32(b, MALI_COMPUTE_SR_JOB_OFFSET_Z), 0); /* Job size */ - cs_move32_to(b, cs_sr_reg32(b, 37), grid.count[0]); - cs_move32_to(b, cs_sr_reg32(b, 38), grid.count[1]); - cs_move32_to(b, cs_sr_reg32(b, 39), grid.count[2]); + cs_move32_to(b, cs_sr_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_X), + grid.count[0]); + cs_move32_to(b, cs_sr_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_Y), + grid.count[1]); + cs_move32_to(b, cs_sr_reg32(b, MALI_COMPUTE_SR_JOB_SIZE_Z), + grid.count[2]); } panvk_per_arch(cs_pick_iter_sb)(cmdbuf, PANVK_SUBQUEUE_COMPUTE);