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

intel/compiler: Restrict cs_threads to 64

Browse source 
Our current GPGPU_WALKER code only supports up to 64 threads.

On HSW we could use up to 70 and TGL up to 112, but only if the walker
is adjusted so the width does not exceed 64. Work to support this is
in progress.

Previous to this change, we might try to downgrade to SIMD8 if the
SIMD16 shader spilled. Since HSW and TGL have the max number of
threads above 64, we would then try to emit an invalid GPGPU walker
command.

Fixes: 932045061b ("i965/cs: Emit compute shader code and upload programs")
Signed-off-by: Jordan Justen <jordan.l.justen@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Paulo Zanoni <paulo.r.zanoni@intel.com>
Tested-by: Paulo Zanoni <paulo.r.zanoni@intel.com>
This commit is contained in:
Jordan Justen 2019-10-24 11:55:23 -07:00
parent 0932363489
commit cf12faef61
No known key found for this signature in database
GPG key ID: 37F99F68CAF992EB
1 changed files with 3 additions and 1 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

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

@ -8841,8 +8841,10 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data,
src_shader->info.cs.local_size[0] * src_shader->info.cs.local_size[1] * src_shader->info.cs.local_size[0] * src_shader->info.cs.local_size[1] *
src_shader->info.cs.local_size[2]; src_shader->info.cs.local_size[2];
/* Limit max_threads to 64 for the GPGPU_WALKER command */
const uint32_t max_threads = MIN2(64, compiler->devinfo->max_cs_threads);
unsigned min_dispatch_width = unsigned min_dispatch_width =
DIV_ROUND_UP(local_workgroup_size, compiler->devinfo->max_cs_threads); DIV_ROUND_UP(local_workgroup_size, max_threads);
min_dispatch_width = MAX2(8, min_dispatch_width); min_dispatch_width = MAX2(8, min_dispatch_width);
min_dispatch_width = util_next_power_of_two(min_dispatch_width); min_dispatch_width = util_next_power_of_two(min_dispatch_width);
assert(min_dispatch_width <= 32); assert(min_dispatch_width <= 32);