From 1cdc4be14b66108ae0e8069686ac3efe52bef3cb Mon Sep 17 00:00:00 2001 From: Caio Oliveira Date: Thu, 21 Sep 2023 13:35:42 -0700 Subject: [PATCH] intel/compiler: Don't allocate memory for SIMD select error handling The position in the error array already indicate the SIMD in question, so take off all the formatted printing from the errors -- which in some cases were just not needed. We lose a little bit of extra context but it is all easily derivable from the message and the SIMD. This also will remove the overhead when SIMD selection is being used to just to find the selected dispatch width -- at a point where the shaders were already compiled -- and the errors are not used at all. Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/9849 Reviewed-by: Kenneth Graunke Part-of: --- src/intel/compiler/brw_fs.cpp | 8 ++--- src/intel/compiler/brw_mesh.cpp | 10 +++---- src/intel/compiler/brw_private.h | 1 - src/intel/compiler/brw_simd_selection.cpp | 35 +++++----------------- src/intel/compiler/test_simd_selection.cpp | 1 - 5 files changed, 17 insertions(+), 38 deletions(-) diff --git a/src/intel/compiler/brw_fs.cpp b/src/intel/compiler/brw_fs.cpp index 0e7960f9472..c67f9044c9f 100644 --- a/src/intel/compiler/brw_fs.cpp +++ b/src/intel/compiler/brw_fs.cpp @@ -8202,7 +8202,6 @@ brw_compile_cs(const struct brw_compiler *compiler, } brw_simd_selection_state simd_state{ - .mem_ctx = params->base.mem_ctx, .devinfo = compiler->devinfo, .prog_data = prog_data, .required_width = brw_required_dispatch_width(&nir->info), @@ -8260,7 +8259,8 @@ brw_compile_cs(const struct brw_compiler *compiler, if (selected_simd < 0) { params->base.error_str = ralloc_asprintf(params->base.mem_ctx, - "Can't compile shader: %s, %s and %s.\n", + "Can't compile shader: " + "SIMD8 '%s', SIMD16 '%s' and SIMD32 '%s'.\n", simd_state.error[0], simd_state.error[1], simd_state.error[2]); return NULL; @@ -8353,7 +8353,6 @@ compile_single_bs(const struct brw_compiler *compiler, key->base.robust_flags); brw_simd_selection_state simd_state{ - .mem_ctx = params->base.mem_ctx, .devinfo = compiler->devinfo, .prog_data = prog_data, @@ -8396,7 +8395,8 @@ compile_single_bs(const struct brw_compiler *compiler, if (selected_simd < 0) { params->base.error_str = ralloc_asprintf(params->base.mem_ctx, - "Can't compile shader: %s and %s.", + "Can't compile shader: " + "SIMD8 '%s' and SIMD16 '%s'.\n", simd_state.error[0], simd_state.error[1]); return 0; } diff --git a/src/intel/compiler/brw_mesh.cpp b/src/intel/compiler/brw_mesh.cpp index e889fad3cc4..1e4b3a55083 100644 --- a/src/intel/compiler/brw_mesh.cpp +++ b/src/intel/compiler/brw_mesh.cpp @@ -298,7 +298,6 @@ brw_compile_task(const struct brw_compiler *compiler, BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_DRAW_ID); brw_simd_selection_state simd_state{ - .mem_ctx = params->base.mem_ctx, .devinfo = compiler->devinfo, .prog_data = &prog_data->base, .required_width = brw_required_dispatch_width(&nir->info), @@ -344,7 +343,8 @@ brw_compile_task(const struct brw_compiler *compiler, if (selected_simd < 0) { params->base.error_str = ralloc_asprintf(params->base.mem_ctx, - "Can't compile shader: %s, %s and %s.\n", + "Can't compile shader: " + "SIMD8 '%s', SIMD16 '%s' and SIMD32 '%s'.\n", simd_state.error[0], simd_state.error[1], simd_state.error[2]); return NULL; @@ -1513,7 +1513,6 @@ brw_compile_mesh(const struct brw_compiler *compiler, brw_nir_lower_mue_outputs(nir, &prog_data->map); brw_simd_selection_state simd_state{ - .mem_ctx = params->base.mem_ctx, .devinfo = compiler->devinfo, .prog_data = &prog_data->base, .required_width = brw_required_dispatch_width(&nir->info), @@ -1571,9 +1570,10 @@ brw_compile_mesh(const struct brw_compiler *compiler, if (selected_simd < 0) { params->base.error_str = ralloc_asprintf(params->base.mem_ctx, - "Can't compile shader: %s, %s and %s.\n", + "Can't compile shader: " + "SIMD8 '%s', SIMD16 '%s' and SIMD32 '%s'.\n", simd_state.error[0], simd_state.error[1], - simd_state.error[2]);; + simd_state.error[2]); return NULL; } diff --git a/src/intel/compiler/brw_private.h b/src/intel/compiler/brw_private.h index a7ebb88acfb..922ec8abc31 100644 --- a/src/intel/compiler/brw_private.h +++ b/src/intel/compiler/brw_private.h @@ -34,7 +34,6 @@ unsigned brw_required_dispatch_width(const struct shader_info *info); static constexpr int SIMD_COUNT = 3; struct brw_simd_selection_state { - void *mem_ctx; const struct intel_device_info *devinfo; std::variant 0 && state.compiled[simd - 1] && workgroup_size <= (width / 2)) { - state.error[simd] = ralloc_asprintf( - state.mem_ctx, "SIMD%u skipped because workgroup size %u already fits in SIMD%u", - width, workgroup_size, width / 2); + state.error[simd] = "Workgroup size already fits in smaller SIMD"; return false; } if (DIV_ROUND_UP(workgroup_size, width) > max_threads) { - state.error[simd] = ralloc_asprintf( - state.mem_ctx, "SIMD%u can't fit all %u invocations in %u threads", - width, workgroup_size, max_threads); + state.error[simd] = "Would need more than max_threads to fit all invocations"; return false; } } @@ -129,24 +122,19 @@ brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd) */ if (width == 32) { if (!INTEL_DEBUG(DEBUG_DO32) && (state.compiled[0] || state.compiled[1])) { - state.error[simd] = ralloc_strdup( - state.mem_ctx, "SIMD32 skipped because not required"); + state.error[simd] = "SIMD32 not required (use INTEL_DEBUG=do32 to force)"; return false; } } } if (width == 32 && cs_prog_data && cs_prog_data->base.ray_queries > 0) { - state.error[simd] = ralloc_asprintf( - state.mem_ctx, "SIMD%u skipped because of ray queries", - width); + state.error[simd] = "Ray queries not supported"; return false; } if (width == 32 && cs_prog_data && cs_prog_data->uses_btd_stack_ids) { - state.error[simd] = ralloc_asprintf( - state.mem_ctx, "SIMD%u skipped because of bindless shader calls", - width); + state.error[simd] = "Bindless shader calls not supported"; return false; } @@ -182,9 +170,7 @@ brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd) static_assert(ARRAY_SIZE(env_skip) == SIMD_COUNT); if (unlikely(env_skip[simd])) { - state.error[simd] = ralloc_asprintf( - state.mem_ctx, "SIMD%u skipped because INTEL_DEBUG=no%u", - width, width); + state.error[simd] = "Disabled by INTEL_DEBUG environment variable"; return false; } @@ -257,10 +243,7 @@ brw_simd_select_for_workgroup_size(const struct intel_device_info *devinfo, cloned.prog_mask = 0; cloned.prog_spilled = 0; - void *mem_ctx = ralloc_context(NULL); - brw_simd_selection_state simd_state{ - .mem_ctx = mem_ctx, .devinfo = devinfo, .prog_data = &cloned, }; @@ -275,7 +258,5 @@ brw_simd_select_for_workgroup_size(const struct intel_device_info *devinfo, } } - ralloc_free(mem_ctx); - return brw_simd_select(simd_state); } diff --git a/src/intel/compiler/test_simd_selection.cpp b/src/intel/compiler/test_simd_selection.cpp index abd3ce070db..ce7490ebc6e 100644 --- a/src/intel/compiler/test_simd_selection.cpp +++ b/src/intel/compiler/test_simd_selection.cpp @@ -46,7 +46,6 @@ protected: , devinfo(rzalloc(mem_ctx, intel_device_info)) , prog_data(rzalloc(mem_ctx, struct brw_cs_prog_data)) , simd_state{ - .mem_ctx = mem_ctx, .devinfo = devinfo, .prog_data = prog_data, }