[Mesa-dev] [PATCH] radeonsi: store group_size_variable in struct si_compute

[Marek Olšák]

Reviewed-by: Marek Olšák <marek.olsak at amd.com>

Marek

On Fri, Nov 18, 2016 at 8:22 PM, Nicolai Hähnle <nhaehnle at gmail.com> wrote:
> From: Nicolai Hähnle <nicolai.haehnle at 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-stable at 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 at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/mesa-dev