summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorJordan Justen <jordan.l.justen@intel.com>2019-10-24 11:55:23 -0700
committerEric Engestrom <eric@engestrom.ch>2020-03-06 22:59:30 +0100
commit8e986a28fabafefd48b182a1e82990dd11536295 (patch)
treed65c399316fe975bb14a77ce3dc7ec90550cbd62
parent0bea36c204b78198ea613adb147527c6fa231ae3 (diff)
intel/compiler: Restrict cs_threads to 64
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: 932045061b5 ("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> (cherry picked from commit cf12faef614ab7cd9996410f1d161558a3853936)
-rw-r--r--.pick_status.json2
-rw-r--r--src/intel/compiler/brw_fs.cpp4
2 files changed, 4 insertions, 2 deletions
diff --git a/.pick_status.json b/.pick_status.json
index 909216f47f8..7df243e9fc9 100644
--- a/.pick_status.json
+++ b/.pick_status.json
@@ -2209,7 +2209,7 @@
"description": "intel/compiler: Restrict cs_threads to 64",
"nominated": true,
"nomination_type": 1,
- "resolution": 0,
+ "resolution": 1,
"master_sha": null,
"because_sha": "932045061b5850368e8a4a5b3e6609eba6ed8d66"
},
diff --git a/src/intel/compiler/brw_fs.cpp b/src/intel/compiler/brw_fs.cpp
index b0abb8c9602..2b4e7a2aa5a 100644
--- a/src/intel/compiler/brw_fs.cpp
+++ b/src/intel/compiler/brw_fs.cpp
@@ -8549,8 +8549,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[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 =
- 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 = util_next_power_of_two(min_dispatch_width);
assert(min_dispatch_width <= 32);