[Mesa-dev] [RFC 11/11] i965/compiler: Take an explicit shader_info parameter in compile_cs
Jason Ekstrand
jason at jlekstrand.net
Fri Oct 28 20:46:40 UTC 2016
---
src/intel/vulkan/anv_pipeline.c | 2 +-
src/mesa/drivers/dri/i965/brw_compiler.h | 1 +
src/mesa/drivers/dri/i965/brw_cs.c | 4 ++--
src/mesa/drivers/dri/i965/brw_fs.cpp | 25 +++++++++++++------------
src/mesa/drivers/dri/i965/brw_nir.h | 2 +-
src/mesa/drivers/dri/i965/brw_nir_intrinsics.c | 9 ++++++---
6 files changed, 24 insertions(+), 19 deletions(-)
diff --git a/src/intel/vulkan/anv_pipeline.c b/src/intel/vulkan/anv_pipeline.c
index dc5a589..639f1c9 100644
--- a/src/intel/vulkan/anv_pipeline.c
+++ b/src/intel/vulkan/anv_pipeline.c
@@ -749,7 +749,7 @@ anv_pipeline_compile_cs(struct anv_pipeline *pipeline,
unsigned code_size;
const unsigned *shader_code =
- brw_compile_cs(compiler, NULL, mem_ctx, &key, &prog_data, nir,
+ brw_compile_cs(compiler, NULL, mem_ctx, &key, &prog_data, nir->info, nir,
-1, &code_size, NULL);
if (shader_code == NULL) {
ralloc_free(mem_ctx);
diff --git a/src/mesa/drivers/dri/i965/brw_compiler.h b/src/mesa/drivers/dri/i965/brw_compiler.h
index 4e306db..d3e99d1 100644
--- a/src/mesa/drivers/dri/i965/brw_compiler.h
+++ b/src/mesa/drivers/dri/i965/brw_compiler.h
@@ -859,6 +859,7 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data,
void *mem_ctx,
const struct brw_cs_prog_key *key,
struct brw_cs_prog_data *prog_data,
+ shader_info *info,
const struct nir_shader *shader,
int shader_time_index,
unsigned *final_assembly_size,
diff --git a/src/mesa/drivers/dri/i965/brw_cs.c b/src/mesa/drivers/dri/i965/brw_cs.c
index d16fff8..87d552c 100644
--- a/src/mesa/drivers/dri/i965/brw_cs.c
+++ b/src/mesa/drivers/dri/i965/brw_cs.c
@@ -125,8 +125,8 @@ brw_codegen_cs_prog(struct brw_context *brw,
char *error_str;
program = brw_compile_cs(brw->screen->compiler, brw, mem_ctx, key,
- &prog_data, cp->program.nir, st_index,
- &program_size, &error_str);
+ &prog_data, &cp->program.info, cp->program.nir,
+ st_index, &program_size, &error_str);
if (program == NULL) {
prog->LinkStatus = false;
ralloc_strcat(&prog->InfoLog, error_str);
diff --git a/src/mesa/drivers/dri/i965/brw_fs.cpp b/src/mesa/drivers/dri/i965/brw_fs.cpp
index 99d535e..f28e766 100644
--- a/src/mesa/drivers/dri/i965/brw_fs.cpp
+++ b/src/mesa/drivers/dri/i965/brw_fs.cpp
@@ -6643,6 +6643,7 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data,
void *mem_ctx,
const struct brw_cs_prog_key *key,
struct brw_cs_prog_data *prog_data,
+ shader_info *info,
const nir_shader *src_shader,
int shader_time_index,
unsigned *final_assembly_size,
@@ -6662,15 +6663,15 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data,
MAX2(shader->num_uniforms,
(unsigned)4 * (prog_data->thread_local_id_index + 1));
- brw_nir_lower_intrinsics(shader, &prog_data->base);
+ brw_nir_lower_intrinsics(shader, info, &prog_data->base);
shader = brw_postprocess_nir(shader, compiler->devinfo, true);
- prog_data->local_size[0] = shader->info->cs.local_size[0];
- prog_data->local_size[1] = shader->info->cs.local_size[1];
- prog_data->local_size[2] = shader->info->cs.local_size[2];
+ prog_data->local_size[0] = info->cs.local_size[0];
+ prog_data->local_size[1] = info->cs.local_size[1];
+ prog_data->local_size[2] = info->cs.local_size[2];
unsigned local_workgroup_size =
- shader->info->cs.local_size[0] * shader->info->cs.local_size[1] *
- shader->info->cs.local_size[2];
+ info->cs.local_size[0] * info->cs.local_size[1] *
+ 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);
@@ -6682,7 +6683,7 @@ 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, shader->info, 8, shader_time_index);
+ shader, info, 8, shader_time_index);
if (simd_required <= 8) {
if (!v8.run_cs()) {
fail_msg = v8.fail_msg;
@@ -6696,7 +6697,7 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data,
fs_visitor v16(compiler, log_data, mem_ctx, key, &prog_data->base,
NULL, /* Never used in core profile */
- shader, shader->info, 16, shader_time_index);
+ shader, info, 16, shader_time_index);
if (likely(!(INTEL_DEBUG & DEBUG_NO16)) &&
!fail_msg && v8.max_dispatch_width >= 16 &&
simd_required <= 16) {
@@ -6722,7 +6723,7 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data,
fs_visitor v32(compiler, log_data, mem_ctx, key, &prog_data->base,
NULL, /* Never used in core profile */
- shader, shader->info, 32, shader_time_index);
+ shader, info, 32, shader_time_index);
if (!fail_msg && v8.max_dispatch_width >= 32 &&
(simd_required > 16 || (INTEL_DEBUG & DEBUG_DO32))) {
/* Try a SIMD32 compile */
@@ -6760,9 +6761,9 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data,
MESA_SHADER_COMPUTE);
if (INTEL_DEBUG & DEBUG_CS) {
char *name = ralloc_asprintf(mem_ctx, "%s compute shader %s",
- shader->info->label ? shader->info->label :
- "unnamed",
- shader->info->name);
+ info->label ? info->label :
+ "unnamed",
+ info->name);
g.enable_debug(name);
}
diff --git a/src/mesa/drivers/dri/i965/brw_nir.h b/src/mesa/drivers/dri/i965/brw_nir.h
index 366d365..2be0459 100644
--- a/src/mesa/drivers/dri/i965/brw_nir.h
+++ b/src/mesa/drivers/dri/i965/brw_nir.h
@@ -95,7 +95,7 @@ void brw_nir_analyze_boolean_resolves(nir_shader *nir);
nir_shader *brw_preprocess_nir(const struct brw_compiler *compiler,
nir_shader *nir);
-bool brw_nir_lower_intrinsics(nir_shader *nir,
+bool brw_nir_lower_intrinsics(nir_shader *nir, const shader_info *info,
struct brw_stage_prog_data *prog_data);
void brw_nir_lower_vs_inputs(nir_shader *nir, shader_info *info,
bool is_scalar,
diff --git a/src/mesa/drivers/dri/i965/brw_nir_intrinsics.c b/src/mesa/drivers/dri/i965/brw_nir_intrinsics.c
index 901a1fb..740a896 100644
--- a/src/mesa/drivers/dri/i965/brw_nir_intrinsics.c
+++ b/src/mesa/drivers/dri/i965/brw_nir_intrinsics.c
@@ -26,6 +26,7 @@
struct lower_intrinsics_state {
nir_shader *nir;
+ const shader_info *info;
union {
struct brw_stage_prog_data *prog_data;
struct brw_cs_prog_data *cs_prog_data;
@@ -41,7 +42,7 @@ read_thread_local_id(struct lower_intrinsics_state *state)
{
nir_builder *b = &state->builder;
nir_shader *nir = state->nir;
- const unsigned *sizes = nir->info->cs.local_size;
+ const unsigned *sizes = state->info->cs.local_size;
const unsigned group_size = sizes[0] * sizes[1] * sizes[2];
/* Some programs have local_size dimensions so small that the thread local
@@ -111,7 +112,7 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state,
* (gl_WorkGroupSize.x * gl_WorkGroupSize.y)) %
* gl_WorkGroupSize.z;
*/
- unsigned *size = nir->info->cs.local_size;
+ const unsigned *size = state->info->cs.local_size;
nir_ssa_def *local_index = nir_load_local_invocation_index(b);
@@ -156,7 +157,8 @@ lower_cs_intrinsics_convert_impl(struct lower_intrinsics_state *state)
}
bool
-brw_nir_lower_intrinsics(nir_shader *nir, struct brw_stage_prog_data *prog_data)
+brw_nir_lower_intrinsics(nir_shader *nir, const shader_info *info,
+ struct brw_stage_prog_data *prog_data)
{
/* Currently we only lower intrinsics for compute shaders */
if (nir->stage != MESA_SHADER_COMPUTE)
@@ -166,6 +168,7 @@ brw_nir_lower_intrinsics(nir_shader *nir, struct brw_stage_prog_data *prog_data)
struct lower_intrinsics_state state;
memset(&state, 0, sizeof(state));
state.nir = nir;
+ state.info = info;
state.prog_data = prog_data;
do {
--
2.5.0.400.gff86faf
More information about the mesa-dev
mailing list