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, }