fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2025-12-30 18:40:13 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

intel/fs/xe2+: Stop building SIMD8 compute-like shaders (CS/BS/TS/MS).

Browse source 
SIMD8 kernels are no longer able to utilize the ALUs efficiently,
since they have twice the vector width as previous platforms.  However
even though there aren't many reasons to use it, SIMD8 is still
supported by the instruction set technically, and it will still be
used for some SIMD-lowering sequences.

Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26605>
This commit is contained in:
Francisco Jerez 2023-12-01 20:35:31 -08:00 • committed by Caio Oliveira
parent 69cc72e50a
commit 7397ba61c2
2 changed files with 11 additions and 2 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

3
src/intel/compiler/brw_fs.cpp
View file

@ -8126,6 +8126,9 @@ compile_single_bs(const struct brw_compiler *compiler,
const unsigned dispatch_width = 8u << simd;
if (dispatch_width == 8 && compiler->devinfo->ver >= 20)
continue;
v[simd] = std::make_unique<fs_visitor>(compiler, &params->base,
&key->base,
&prog_data->base, shader,

10
src/intel/compiler/brw_simd_selection.cpp
View file

@ -104,7 +104,8 @@ brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd)
unsigned max_threads = state.devinfo->max_cs_workgroup_threads;
if (simd > 0 && state.compiled[simd - 1] &&
const unsigned min_simd = state.devinfo->ver >= 20 ? 1 : 0;
if (simd > min_simd && state.compiled[simd - 1] &&
workgroup_size <= (width / 2)) {
state.error[simd] = "Workgroup size already fits in smaller SIMD";
return false;
@ -120,7 +121,7 @@ brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd)
*
* TODO: Use performance_analysis and drop this rule.
*/
if (width == 32) {
if (width == 32 && state.devinfo->ver < 20) {
if (!INTEL_DEBUG(DEBUG_DO32) && (state.compiled[0] || state.compiled[1])) {
state.error[simd] = "SIMD32 not required (use INTEL_DEBUG=do32 to force)";
return false;
@ -128,6 +129,11 @@ brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd)
}
}
if (width == 8 && state.devinfo->ver >= 20) {
state.error[simd] = "SIMD8 not supported on Xe2+";
return false;
}
if (width == 32 && cs_prog_data && cs_prog_data->base.ray_queries > 0) {
state.error[simd] = "Ray queries not supported";
return false;