Reviewed-by: Marek Olšák <marek.ol...@amd.com>
Marek
On Fri, Nov 18, 2016 at 8:22 PM, Nicolai Hähnle <nhaeh...@gmail.com> wrote:
> From: Nicolai Hähnle <nicolai.haeh...@amd.com>
>
> For compute shaders, we free the selector after the shader has been
> compiled, so we need to save this bit somewhere else. Also, make sure that
> this type of bug cannot re-appear, by NULL-ing the selector pointer after
> we're done with it.
>
> This bug has been there since the feature was added, but was only exposed
> in piglit arb_compute_variable_group_size-local-size by commit
> 9bfee7047b70cb0aa026ca9536465762f96cb2b1 (which is totally unrelated).
>
> Cc: 13.0 <mesa-sta...@lists.freedesktop.org>
> ---
> src/gallium/drivers/radeonsi/si_compute.c | 13 ++++++++-----
> 1 file changed, 8 insertions(+), 5 deletions(-)
>
> diff --git a/src/gallium/drivers/radeonsi/si_compute.c
> b/src/gallium/drivers/radeonsi/si_compute.c
> index f1887bb..69d57b9 100644
> --- a/src/gallium/drivers/radeonsi/si_compute.c
> +++ b/src/gallium/drivers/radeonsi/si_compute.c
> @@ -35,21 +35,22 @@
> #define MAX_GLOBAL_BUFFERS 20
>
> struct si_compute {
> unsigned ir_type;
> unsigned local_size;
> unsigned private_size;
> unsigned input_size;
> struct si_shader shader;
>
> struct pipe_resource *global_buffers[MAX_GLOBAL_BUFFERS];
> - bool use_code_object_v2;
> + unsigned use_code_object_v2 : 1;
> + unsigned variable_group_size : 1;
> };
>
> struct dispatch_packet {
> uint16_t header;
> uint16_t setup;
> uint16_t workgroup_size_x;
> uint16_t workgroup_size_y;
> uint16_t workgroup_size_z;
> uint16_t reserved0;
> uint32_t grid_size_x;
> @@ -140,21 +141,25 @@ static void *si_create_compute_state(
> S_00B848_SGPRS((shader->config.num_sgprs - 1) / 8)
> |
> S_00B848_DX10_CLAMP(1) |
> S_00B848_FLOAT_MODE(shader->config.float_mode);
>
> shader->config.rsrc2 =
> S_00B84C_USER_SGPR(SI_CS_NUM_USER_SGPR) |
> S_00B84C_SCRATCH_EN(scratch_enabled) |
> S_00B84C_TGID_X_EN(1) | S_00B84C_TGID_Y_EN(1) |
> S_00B84C_TGID_Z_EN(1) | S_00B84C_TIDIG_COMP_CNT(2)
> |
> S_00B84C_LDS_SIZE(shader->config.lds_size);
>
> + program->variable_group_size =
> +
> sel.info.properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] == 0;
> +
> FREE(sel.tokens);
> + program->shader.selector = NULL;
> } else {
> const struct pipe_llvm_program_header *header;
> const char *code;
> header = cso->prog;
> code = cso->prog + sizeof(struct pipe_llvm_program_header);
>
> radeon_elf_read(code, header->num_bytes,
> &program->shader.binary);
> if (program->use_code_object_v2) {
> const amd_kernel_code_t *code_object =
> si_compute_get_code_object(program, 0);
> @@ -600,28 +605,26 @@ static void si_setup_tgsi_grid(struct si_context *sctx,
> radeon_emit(cs, PKT3(PKT3_COPY_DATA, 4, 0));
> radeon_emit(cs, COPY_DATA_SRC_SEL(COPY_DATA_MEM) |
> COPY_DATA_DST_SEL(COPY_DATA_REG));
> radeon_emit(cs, (va + 4 * i));
> radeon_emit(cs, (va + 4 * i) >> 32);
> radeon_emit(cs, (grid_size_reg >> 2) + i);
> radeon_emit(cs, 0);
> }
> } else {
> struct si_compute *program = sctx->cs_shader_state.program;
> - bool variable_group_size =
> -
> program->shader.selector->info.properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH]
> == 0;
>
> - radeon_set_sh_reg_seq(cs, grid_size_reg, variable_group_size
> ? 6 : 3);
> + radeon_set_sh_reg_seq(cs, grid_size_reg,
> program->variable_group_size ? 6 : 3);
> radeon_emit(cs, info->grid[0]);
> radeon_emit(cs, info->grid[1]);
> radeon_emit(cs, info->grid[2]);
> - if (variable_group_size) {
> + if (program->variable_group_size) {
> radeon_emit(cs, info->block[0]);
> radeon_emit(cs, info->block[1]);
> radeon_emit(cs, info->block[2]);
> }
> }
> }
>
> static void si_emit_dispatch_packets(struct si_context *sctx,
> const struct pipe_grid_info *info)
> {
> --
> 2.7.4
>
> _______________________________________________
> mesa-dev mailing list
> mesa-dev@lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/mesa-dev
_______________________________________________
mesa-dev mailing list
mesa-dev@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/mesa-dev
