From 9377dc417dfc85cbf60b3e0b7ae113452814c1d4 Mon Sep 17 00:00:00 2001 From: Caio Oliveira Date: Mon, 22 Jan 2024 23:11:13 -0800 Subject: [PATCH] intel/blorp: Use a struct to return blorp_compile_*() results Allocate the prog_data instead of expecting one from the user, also explicitly return both kernel and prog_data size, so that the plumbing code isn't required to use the exact prog_data type. Reviewed-by: Kenneth Graunke Part-of: --- src/intel/blorp/blorp.c | 48 ++++++++++++++++++++++------------- src/intel/blorp/blorp_blit.c | 21 ++++++--------- src/intel/blorp/blorp_clear.c | 41 ++++++++++++------------------ src/intel/blorp/blorp_priv.h | 23 ++++++++++------- 4 files changed, 69 insertions(+), 64 deletions(-) diff --git a/src/intel/blorp/blorp.c b/src/intel/blorp/blorp.c index c88ddc90b9b..b6107dfca4e 100644 --- a/src/intel/blorp/blorp.c +++ b/src/intel/blorp/blorp.c @@ -250,18 +250,16 @@ blorp_params_init(struct blorp_params *params) params->num_layers = 1; } -const unsigned * +struct blorp_program blorp_compile_fs(struct blorp_context *blorp, void *mem_ctx, struct nir_shader *nir, bool multisample_fbo, - bool use_repclear, - struct brw_wm_prog_data *wm_prog_data) + bool use_repclear) { const struct brw_compiler *compiler = blorp->compiler; nir->options = compiler->nir_options[MESA_SHADER_FRAGMENT]; - memset(wm_prog_data, 0, sizeof(*wm_prog_data)); - + struct brw_wm_prog_data *wm_prog_data = rzalloc(mem_ctx, struct brw_wm_prog_data); wm_prog_data->base.nr_params = 0; wm_prog_data->base.param = NULL; @@ -296,13 +294,18 @@ blorp_compile_fs(struct blorp_context *blorp, void *mem_ctx, .max_polygons = 1, }; - return brw_compile_fs(compiler, ¶ms); + const unsigned *kernel = brw_compile_fs(compiler, ¶ms); + return (struct blorp_program){ + .kernel = kernel, + .kernel_size = wm_prog_data->base.program_size, + .prog_data = wm_prog_data, + .prog_data_size = sizeof(*wm_prog_data), + }; } -const unsigned * +struct blorp_program blorp_compile_vs(struct blorp_context *blorp, void *mem_ctx, - struct nir_shader *nir, - struct brw_vs_prog_data *vs_prog_data) + struct nir_shader *nir) { const struct brw_compiler *compiler = blorp->compiler; @@ -312,6 +315,7 @@ blorp_compile_vs(struct blorp_context *blorp, void *mem_ctx, brw_preprocess_nir(compiler, nir, &opts); nir_shader_gather_info(nir, nir_shader_get_entrypoint(nir)); + struct brw_vs_prog_data *vs_prog_data = rzalloc(mem_ctx, struct brw_vs_prog_data); vs_prog_data->inputs_read = nir->info.inputs_read; brw_compute_vue_map(compiler->devinfo, @@ -333,7 +337,13 @@ blorp_compile_vs(struct blorp_context *blorp, void *mem_ctx, .prog_data = vs_prog_data, }; - return brw_compile_vs(compiler, ¶ms); + const unsigned *kernel = brw_compile_vs(compiler, ¶ms); + return (struct blorp_program) { + .kernel = kernel, + .kernel_size = vs_prog_data->base.base.program_size, + .prog_data = vs_prog_data, + .prog_data_size = sizeof(*vs_prog_data), + }; } static bool @@ -348,17 +358,14 @@ lower_base_workgroup_id(nir_builder *b, nir_intrinsic_instr *intrin, return true; } -const unsigned * +struct blorp_program blorp_compile_cs(struct blorp_context *blorp, void *mem_ctx, - struct nir_shader *nir, - struct brw_cs_prog_data *cs_prog_data) + struct nir_shader *nir) { const struct brw_compiler *compiler = blorp->compiler; nir->options = compiler->nir_options[MESA_SHADER_COMPUTE]; - memset(cs_prog_data, 0, sizeof(*cs_prog_data)); - struct brw_nir_compiler_opts opts = {}; brw_preprocess_nir(compiler, nir, &opts); nir_shader_gather_info(nir, nir_shader_get_entrypoint(nir)); @@ -370,6 +377,8 @@ blorp_compile_cs(struct blorp_context *blorp, void *mem_ctx, sizeof(struct blorp_wm_inputs)); nir->num_uniforms = offsetof(struct blorp_wm_inputs, subgroup_id); unsigned nr_params = nir->num_uniforms / 4; + + struct brw_cs_prog_data *cs_prog_data = rzalloc(mem_ctx, struct brw_cs_prog_data); cs_prog_data->base.nr_params = nr_params; cs_prog_data->base.param = rzalloc_array(NULL, uint32_t, nr_params); @@ -392,12 +401,17 @@ blorp_compile_cs(struct blorp_context *blorp, void *mem_ctx, .prog_data = cs_prog_data, }; - const unsigned *program = brw_compile_cs(compiler, ¶ms); + const unsigned *kernel = brw_compile_cs(compiler, ¶ms); ralloc_free(cs_prog_data->base.param); cs_prog_data->base.param = NULL; - return program; + return (struct blorp_program) { + .kernel = kernel, + .kernel_size = cs_prog_data->base.program_size, + .prog_data = cs_prog_data, + .prog_data_size = sizeof(*cs_prog_data), + }; } struct blorp_sf_key { diff --git a/src/intel/blorp/blorp_blit.c b/src/intel/blorp/blorp_blit.c index 4b4224579c8..fa426ec8385 100644 --- a/src/intel/blorp/blorp_blit.c +++ b/src/intel/blorp/blorp_blit.c @@ -1510,23 +1510,20 @@ blorp_get_blit_kernel_fs(struct blorp_batch *batch, void *mem_ctx = ralloc_context(NULL); - const unsigned *program; - struct brw_wm_prog_data prog_data; - nir_shader *nir = blorp_build_nir_shader(blorp, batch, mem_ctx, key); nir->info.name = ralloc_strdup(nir, blorp_shader_type_to_name(key->base.shader_type)); const bool multisample_fbo = key->rt_samples > 1; - program = blorp_compile_fs(blorp, mem_ctx, nir, multisample_fbo, false, - &prog_data); + const struct blorp_program p = + blorp_compile_fs(blorp, mem_ctx, nir, multisample_fbo, false); bool result = blorp->upload_shader(batch, MESA_SHADER_FRAGMENT, key, sizeof(*key), - program, prog_data.base.program_size, - &prog_data.base, sizeof(prog_data), + p.kernel, p.kernel_size, + p.prog_data, p.prog_data_size, ¶ms->wm_prog_kernel, ¶ms->wm_prog_data); ralloc_free(mem_ctx); @@ -1546,9 +1543,6 @@ blorp_get_blit_kernel_cs(struct blorp_batch *batch, void *mem_ctx = ralloc_context(NULL); - const unsigned *program; - struct brw_cs_prog_data prog_data; - nir_shader *nir = blorp_build_nir_shader(blorp, batch, mem_ctx, prog_key); nir->info.name = ralloc_strdup(nir, "BLORP-gpgpu-blit"); @@ -1556,13 +1550,14 @@ blorp_get_blit_kernel_cs(struct blorp_batch *batch, assert(prog_key->rt_samples == 1); - program = blorp_compile_cs(blorp, mem_ctx, nir, &prog_data); + const struct blorp_program p = + blorp_compile_cs(blorp, mem_ctx, nir); bool result = blorp->upload_shader(batch, MESA_SHADER_COMPUTE, prog_key, sizeof(*prog_key), - program, prog_data.base.program_size, - &prog_data.base, sizeof(prog_data), + p.kernel, p.kernel_size, + p.prog_data, p.prog_data_size, ¶ms->cs_prog_kernel, ¶ms->cs_prog_data); ralloc_free(mem_ctx); diff --git a/src/intel/blorp/blorp_clear.c b/src/intel/blorp/blorp_clear.c index c0e6d19c600..de784e066e5 100644 --- a/src/intel/blorp/blorp_clear.c +++ b/src/intel/blorp/blorp_clear.c @@ -28,7 +28,6 @@ #include "util/format_srgb.h" #include "blorp_priv.h" -#include "compiler/brw_eu_defines.h" #include "dev/intel_debug.h" #include "blorp_nir_builder.h" @@ -90,17 +89,15 @@ blorp_params_get_clear_kernel_fs(struct blorp_batch *batch, frag_color->data.location = FRAG_RESULT_COLOR; nir_store_var(&b, frag_color, color, 0xf); - struct brw_wm_prog_data prog_data; const bool multisample_fbo = false; - const unsigned *program = - blorp_compile_fs(blorp, mem_ctx, b.shader, multisample_fbo, use_replicated_data, - &prog_data); + struct blorp_program p = + blorp_compile_fs(blorp, mem_ctx, b.shader, multisample_fbo, use_replicated_data); bool result = blorp->upload_shader(batch, MESA_SHADER_FRAGMENT, &blorp_key, sizeof(blorp_key), - program, prog_data.base.program_size, - &prog_data.base, sizeof(prog_data), + p.kernel, p.kernel_size, + p.prog_data, p.prog_data_size, ¶ms->wm_prog_kernel, ¶ms->wm_prog_data); ralloc_free(mem_ctx); @@ -164,15 +161,14 @@ blorp_params_get_clear_kernel_cs(struct blorp_batch *batch, nir_pop_if(&b, NULL); - struct brw_cs_prog_data prog_data; - const unsigned *program = - blorp_compile_cs(blorp, mem_ctx, b.shader, &prog_data); + const struct blorp_program p = + blorp_compile_cs(blorp, mem_ctx, b.shader); bool result = blorp->upload_shader(batch, MESA_SHADER_COMPUTE, &blorp_key, sizeof(blorp_key), - program, prog_data.base.program_size, - &prog_data.base, sizeof(prog_data), + p.kernel, p.kernel_size, + p.prog_data, p.prog_data_size, ¶ms->cs_prog_kernel, ¶ms->cs_prog_data); ralloc_free(mem_ctx); @@ -272,17 +268,14 @@ blorp_params_get_layer_offset_vs(struct blorp_batch *batch, nir_copy_var(&b, v_out, a_in); } - struct brw_vs_prog_data vs_prog_data; - memset(&vs_prog_data, 0, sizeof(vs_prog_data)); - - const unsigned *program = - blorp_compile_vs(blorp, mem_ctx, b.shader, &vs_prog_data); + const struct blorp_program p = + blorp_compile_vs(blorp, mem_ctx, b.shader); bool result = blorp->upload_shader(batch, MESA_SHADER_VERTEX, &blorp_key, sizeof(blorp_key), - program, vs_prog_data.base.base.program_size, - &vs_prog_data.base.base, sizeof(vs_prog_data), + p.kernel, p.kernel_size, + p.prog_data, p.prog_data_size, ¶ms->vs_prog_kernel, ¶ms->vs_prog_data); ralloc_free(mem_ctx); @@ -1381,17 +1374,15 @@ blorp_params_get_mcs_partial_resolve_kernel(struct blorp_batch *batch, } nir_store_var(&b, frag_color, clear_color, 0xf); - struct brw_wm_prog_data prog_data; const bool multisample_fbo = true; - const unsigned *program = - blorp_compile_fs(blorp, mem_ctx, b.shader, multisample_fbo, false, - &prog_data); + const struct blorp_program p = + blorp_compile_fs(blorp, mem_ctx, b.shader, multisample_fbo, false); bool result = blorp->upload_shader(batch, MESA_SHADER_FRAGMENT, &blorp_key, sizeof(blorp_key), - program, prog_data.base.program_size, - &prog_data.base, sizeof(prog_data), + p.kernel, p.kernel_size, + p.prog_data, p.prog_data_size, ¶ms->wm_prog_kernel, ¶ms->wm_prog_data); ralloc_free(mem_ctx); diff --git a/src/intel/blorp/blorp_priv.h b/src/intel/blorp/blorp_priv.h index a9210e2face..2be9b17e311 100644 --- a/src/intel/blorp/blorp_priv.h +++ b/src/intel/blorp/blorp_priv.h @@ -418,17 +418,23 @@ bool blorp_blitter_supports_aux(const struct intel_device_info *devinfo, const char *blorp_shader_type_to_name(enum blorp_shader_type type); const char *blorp_shader_pipeline_to_name(enum blorp_shader_pipeline pipe); -const unsigned * +struct blorp_program { + const void *kernel; + uint32_t kernel_size; + + const void *prog_data; + uint32_t prog_data_size; +}; + +struct blorp_program blorp_compile_fs(struct blorp_context *blorp, void *mem_ctx, struct nir_shader *nir, bool multisample_fbo, - bool use_repclear, - struct brw_wm_prog_data *wm_prog_data); + bool use_repclear); -const unsigned * +struct blorp_program blorp_compile_vs(struct blorp_context *blorp, void *mem_ctx, - struct nir_shader *nir, - struct brw_vs_prog_data *vs_prog_data); + struct nir_shader *nir); bool blorp_ensure_sf_program(struct blorp_batch *batch, @@ -457,10 +463,9 @@ blorp_set_cs_dims(struct nir_shader *nir, uint8_t local_y) nir->info.workgroup_size[2] = 1; } -const unsigned * +struct blorp_program blorp_compile_cs(struct blorp_context *blorp, void *mem_ctx, - struct nir_shader *nir, - struct brw_cs_prog_data *cs_prog_data); + struct nir_shader *nir); /** \} */