From 4e711872d024ce41c8b07b1150d8a393de21e26d Mon Sep 17 00:00:00 2001 From: Jason Ekstrand Date: Thu, 8 Oct 2015 15:28:26 -0700 Subject: [PATCH] i965/cs: Rework cs_emit to take a nir_shader and a brw_compiler This commit removes all dependence on GL state by getting rid of the brw_context parameter and the GL data structures. Reviewed-by: Topi Pohjolainen --- src/mesa/drivers/dri/i965/brw_cs.c | 10 ++++-- src/mesa/drivers/dri/i965/brw_cs.h | 10 +++--- src/mesa/drivers/dri/i965/brw_fs.cpp | 51 +++++++++++++++------------- 3 files changed, 42 insertions(+), 29 deletions(-) diff --git a/src/mesa/drivers/dri/i965/brw_cs.c b/src/mesa/drivers/dri/i965/brw_cs.c index 45fb816c160..12e75097f8b 100644 --- a/src/mesa/drivers/dri/i965/brw_cs.c +++ b/src/mesa/drivers/dri/i965/brw_cs.c @@ -105,9 +105,15 @@ brw_codegen_cs_prog(struct brw_context *brw, if (INTEL_DEBUG & DEBUG_SHADER_TIME) st_index = brw_get_shader_time_index(brw, prog, &cp->program.Base, ST_CS); - program = brw_cs_emit(brw, mem_ctx, key, &prog_data, - &cp->program, prog, st_index, &program_size); + char *error_str; + program = brw_cs_emit(brw->intelScreen->compiler, brw, mem_ctx, + key, &prog_data, cp->program.Base.nir, + st_index, &program_size, &error_str); if (program == NULL) { + prog->LinkStatus = false; + ralloc_strcat(&prog->InfoLog, error_str); + _mesa_problem(NULL, "Failed to compile compute shader: %s\n", error_str); + ralloc_free(mem_ctx); return false; } diff --git a/src/mesa/drivers/dri/i965/brw_cs.h b/src/mesa/drivers/dri/i965/brw_cs.h index 17c2ff9871a..1a9613e3039 100644 --- a/src/mesa/drivers/dri/i965/brw_cs.h +++ b/src/mesa/drivers/dri/i965/brw_cs.h @@ -39,15 +39,17 @@ extern "C" { void brw_upload_cs_prog(struct brw_context *brw); +struct nir_shader; + const unsigned * -brw_cs_emit(struct brw_context *brw, +brw_cs_emit(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, - struct gl_compute_program *cp, - struct gl_shader_program *prog, + const struct nir_shader *shader, int shader_time_index, - unsigned *final_assembly_size); + unsigned *final_assembly_size, + char **error_str); void brw_cs_fill_local_id_payload(const struct brw_cs_prog_data *cs_prog_data, diff --git a/src/mesa/drivers/dri/i965/brw_fs.cpp b/src/mesa/drivers/dri/i965/brw_fs.cpp index d37a9ed0b55..ce130dffad6 100644 --- a/src/mesa/drivers/dri/i965/brw_fs.cpp +++ b/src/mesa/drivers/dri/i965/brw_fs.cpp @@ -5234,29 +5234,32 @@ fs_visitor::emit_cs_work_group_id_setup() } const unsigned * -brw_cs_emit(struct brw_context *brw, +brw_cs_emit(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, - struct gl_compute_program *cp, - struct gl_shader_program *prog, + const nir_shader *shader, int shader_time_index, - unsigned *final_assembly_size) + unsigned *final_assembly_size, + char **error_str) { - prog_data->local_size[0] = cp->LocalSize[0]; - prog_data->local_size[1] = cp->LocalSize[1]; - prog_data->local_size[2] = cp->LocalSize[2]; + 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]; unsigned local_workgroup_size = - cp->LocalSize[0] * cp->LocalSize[1] * cp->LocalSize[2]; - unsigned max_cs_threads = brw->intelScreen->compiler->devinfo->max_cs_threads; + shader->info.cs.local_size[0] * shader->info.cs.local_size[1] * + shader->info.cs.local_size[2]; + + unsigned max_cs_threads = compiler->devinfo->max_cs_threads; cfg_t *cfg = NULL; const char *fail_msg = NULL; /* Now the main event: Visit the shader IR and generate our CS IR for it. */ - fs_visitor v8(brw->intelScreen->compiler, brw, mem_ctx, key, - &prog_data->base, &cp->Base, cp->Base.nir, 8, shader_time_index); + 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) { @@ -5264,15 +5267,18 @@ brw_cs_emit(struct brw_context *brw, prog_data->simd_size = 8; } - fs_visitor v16(brw->intelScreen->compiler, brw, mem_ctx, key, - &prog_data->base, &cp->Base, cp->Base.nir, 16, shader_time_index); + fs_visitor v16(compiler, log_data, mem_ctx, key, &prog_data->base, + NULL, /* Never used in core profile */ + shader, 16, shader_time_index); if (likely(!(INTEL_DEBUG & DEBUG_NO16)) && !fail_msg && !v8.simd16_unsupported && local_workgroup_size <= 16 * max_cs_threads) { /* Try a SIMD16 compile */ v16.import_uniforms(&v8); if (!v16.run_cs()) { - perf_debug("SIMD16 shader failed to compile: %s", v16.fail_msg); + compiler->shader_perf_log(log_data, + "SIMD16 shader failed to compile: %s", + v16.fail_msg); if (!cfg) { fail_msg = "Couldn't generate SIMD16 program and not " @@ -5286,20 +5292,19 @@ brw_cs_emit(struct brw_context *brw, if (unlikely(cfg == NULL)) { assert(fail_msg); - prog->LinkStatus = false; - ralloc_strcat(&prog->InfoLog, fail_msg); - _mesa_problem(NULL, "Failed to compile compute shader: %s\n", - fail_msg); + if (error_str) + *error_str = ralloc_strdup(mem_ctx, fail_msg); + return NULL; } - fs_generator g(brw->intelScreen->compiler, brw, - mem_ctx, (void*) key, &prog_data->base, + fs_generator g(compiler, log_data, mem_ctx, (void*) key, &prog_data->base, v8.promoted_constants, v8.runtime_check_aads_emit, "CS"); if (INTEL_DEBUG & DEBUG_CS) { - char *name = ralloc_asprintf(mem_ctx, "%s compute shader %d", - prog->Label ? prog->Label : "unnamed", - prog->Name); + char *name = ralloc_asprintf(mem_ctx, "%s compute shader %s", + shader->info.label ? shader->info.label : + "unnamed", + shader->info.name); g.enable_debug(name); } -- 2.30.2