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;
_______________________________________________
mesa-dev mailing list
mesa-dev@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/mesa-dev
