diff options
Diffstat (limited to 'src')
| -rw-r--r-- | src/mesa/drivers/dri/i965/brw_fs.cpp | 20 |
1 files changed, 12 insertions, 8 deletions
diff --git a/src/mesa/drivers/dri/i965/brw_fs.cpp b/src/mesa/drivers/dri/i965/brw_fs.cpp index 27f2123b7a0..b5404e1df50 100644 --- a/src/mesa/drivers/dri/i965/brw_fs.cpp +++ b/src/mesa/drivers/dri/i965/brw_fs.cpp @@ -1933,8 +1933,8 @@ fs_visitor::compact_virtual_grfs() void fs_visitor::assign_constant_locations() { - /* Only the first compile (SIMD8 mode) gets to decide on locations. */ - if (dispatch_width != 8) + /* Only the first compile gets to decide on locations. */ + if (dispatch_width != min_dispatch_width) return; unsigned int num_pull_constants = 0; @@ -5733,6 +5733,7 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data, shader->info.cs.local_size[2]; unsigned max_cs_threads = compiler->devinfo->max_cs_threads; + unsigned simd_required = DIV_ROUND_UP(local_workgroup_size, max_cs_threads); cfg_t *cfg = NULL; const char *fail_msg = NULL; @@ -5742,11 +5743,13 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data, fs_visitor v8(compiler, log_data, mem_ctx, key, &prog_data->base, NULL, /* Never used in core profile */ shader, 8, shader_time_index); - if (!v8.run_cs()) { - fail_msg = v8.fail_msg; - } else if (local_workgroup_size <= 8 * max_cs_threads) { - cfg = v8.cfg; - prog_data->simd_size = 8; + if (simd_required <= 8) { + if (!v8.run_cs()) { + fail_msg = v8.fail_msg; + } else { + cfg = v8.cfg; + prog_data->simd_size = 8; + } } fs_visitor v16(compiler, log_data, mem_ctx, key, &prog_data->base, @@ -5756,7 +5759,8 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data, !fail_msg && !v8.simd16_unsupported && local_workgroup_size <= 16 * max_cs_threads) { /* Try a SIMD16 compile */ - v16.import_uniforms(&v8); + if (simd_required <= 8) + v16.import_uniforms(&v8); if (!v16.run_cs()) { compiler->shader_perf_log(log_data, "SIMD16 shader failed to compile: %s", |