[Mesa-dev] [PATCH v3 30/48] intel/cs: Re-run final NIR optimizations for each SIMD size

[Iago Toral]

On Wed, 2017-10-25 at 16:26 -0700, Jason Ekstrand wrote:
> With the advent of SPIR-V subgroup operations, compute shaders will
> have
> to be slightly different depending on the SIMD size at which they
> execute.  In order to allow us to do dispatch-width specific things
> in
> NIR, we re-run the final NIR stages for each sIMD width.
> 
> One side-effect of this change is that we start rallocing fs_visitors
> which means we need DECLARE_RALLOC_CXX_OPERATORS.
> ---
>  src/intel/compiler/brw_fs.cpp | 103 ++++++++++++++++++++++++++----
> ------------
>  src/intel/compiler/brw_fs.h   |   2 +
>  2 files changed, 66 insertions(+), 39 deletions(-)
> 
> diff --git a/src/intel/compiler/brw_fs.cpp
> b/src/intel/compiler/brw_fs.cpp
> index c0d4c05..c054537 100644
> --- a/src/intel/compiler/brw_fs.cpp
> +++ b/src/intel/compiler/brw_fs.cpp
> @@ -6770,6 +6770,20 @@ cs_set_simd_size(struct brw_cs_prog_data
> *cs_prog_data, unsigned size)
>     cs_prog_data->threads = (group_size + size - 1) / size;
>  }
>  
> +static nir_shader *
> +compile_cs_to_nir(const struct brw_compiler *compiler,
> +                  void *mem_ctx,
> +                  const struct brw_cs_prog_key *key,
> +                  struct brw_cs_prog_data *prog_data,
> +                  const nir_shader *src_shader,
> +                  unsigned dispatch_width)
> +{
> +   nir_shader *shader = nir_shader_clone(mem_ctx, src_shader);
> +   shader = brw_nir_apply_sampler_key(shader, compiler, &key->tex,
> true);
> +   brw_nir_lower_cs_intrinsics(shader);
> +   return brw_postprocess_nir(shader, compiler, true);
> +}
> +
>  const unsigned *
>  brw_compile_cs(const struct brw_compiler *compiler, void *log_data,
>                 void *mem_ctx,
> @@ -6780,17 +6794,12 @@ brw_compile_cs(const struct brw_compiler
> *compiler, void *log_data,
>                 unsigned *final_assembly_size,
>                 char **error_str)
>  {
> -   nir_shader *shader = nir_shader_clone(mem_ctx, src_shader);
> -   shader = brw_nir_apply_sampler_key(shader, compiler, &key->tex,
> true);
> -   brw_nir_lower_cs_intrinsics(shader);
> -   shader = brw_postprocess_nir(shader, compiler, true);
> -
> -   prog_data->local_size[0] = shader->info.cs.local_size[0];
> -   prog_data->local_size[1] = shader->info.cs.local_size[1];
> -   prog_data->local_size[2] = shader->info.cs.local_size[2];
> +   prog_data->local_size[0] = src_shader->info.cs.local_size[0];
> +   prog_data->local_size[1] = src_shader->info.cs.local_size[1];
> +   prog_data->local_size[2] = src_shader->info.cs.local_size[2];
>     unsigned local_workgroup_size =
> -      shader->info.cs.local_size[0] * shader->info.cs.local_size[1]
> *
> -      shader->info.cs.local_size[2];
> +      src_shader->info.cs.local_size[0] * src_shader-
> >info.cs.local_size[1] *
> +      src_shader->info.cs.local_size[2];
>  
>     unsigned min_dispatch_width =
>        DIV_ROUND_UP(local_workgroup_size, compiler->devinfo-
> >max_cs_threads);
> @@ -6798,71 +6807,87 @@ brw_compile_cs(const struct brw_compiler
> *compiler, void *log_data,
>     min_dispatch_width = util_next_power_of_two(min_dispatch_width);
>     assert(min_dispatch_width <= 32);
>  
> +

Extra blank line

> +   fs_visitor *v8 = NULL, *v16 = NULL, *v32 = NULL;
>     cfg_t *cfg = NULL;
>     const char *fail_msg = NULL;
> +   unsigned promoted_constants;