int
glcpp_preprocess(void *ralloc_ctx, const char **shader, char **info_log,
glcpp_extension_iterator extensions, void *state,
struct gl_context *gl_ctx)
{
int errors;
glcpp_parser_t *parser =
glcpp_parser_create(&gl_ctx->Extensions, extensions, state, gl_ctx->API);
if (! gl_ctx->Const.DisableGLSLLineContinuations)
*shader = remove_line_continuations(parser, *shader);
glcpp_lex_set_source_string (parser, *shader);
glcpp_parser_parse (parser);
if (parser->skip_stack)
glcpp_error (&parser->skip_stack->loc, parser, "Unterminated #if\n");
glcpp_parser_resolve_implicit_version(parser);
ralloc_strcat(info_log, parser->info_log->buf);
/* Crimp the buffer first, to conserve memory */
_mesa_string_buffer_crimp_to_fit(parser->output);
ralloc_steal(ralloc_ctx, parser->output->buf);
*shader = parser->output->buf;
errors = parser->error;
glcpp_parser_destroy (parser);
return errors;
}
virtual ir_visitor_status visit_leave(ir_dereference_array *ir)
{
ir_constant *index = ir->array_index->as_constant();
int i;
if (index) {
i = index->value.i[0];
} else {
/* GLSL 1.10 and 1.20 allowed variable sampler array indices,
* while GLSL 1.30 requires that the array indices be
* constant integer expressions. We don't expect any driver
* to actually work with a really variable array index, so
* all that would work would be an unrolled loop counter that ends
* up being constant above.
*/
ralloc_strcat(&shader_program->InfoLog,
"warning: Variable sampler array index unsupported.\n"
"This feature of the language was removed in GLSL 1.20 "
"and is unlikely to be supported for 1.10 in Mesa.\n");
i = 0;
}
if (ir != last) {
this->name = ralloc_asprintf(mem_ctx, "%s[%d]", name, i);
} else {
offset = i;
}
return visit_continue;
}
extern "C" const unsigned *
brw_gs_emit(struct brw_context *brw,
struct gl_shader_program *prog,
struct brw_gs_compile *c,
void *mem_ctx,
unsigned *final_assembly_size)
{
struct brw_shader *shader =
(brw_shader *) prog->_LinkedShaders[MESA_SHADER_GEOMETRY];
if (unlikely(INTEL_DEBUG & DEBUG_GS)) {
printf("GLSL IR for native geometry shader %d:\n", prog->Name);
_mesa_print_ir(shader->ir, NULL);
printf("\n\n");
}
vec4_gs_visitor v(brw, c, prog, shader, mem_ctx);
if (!v.run()) {
prog->LinkStatus = false;
ralloc_strcat(&prog->InfoLog, v.fail_msg);
return NULL;
}
vec4_generator g(brw, prog, &c->gp->program.Base, &c->prog_data.base,
mem_ctx, INTEL_DEBUG & DEBUG_GS);
const unsigned *generated =
g.generate_assembly(&v.instructions, final_assembly_size);
return generated;
}
static void
shader_error(struct gl_context *ctx, struct gl_program *prog, const char *msg)
{
struct gl_shader_program *shader;
shader = _mesa_lookup_shader_program(ctx, prog->Id);
if (shader) {
ralloc_strcat(&shader->InfoLog, msg);
shader->LinkStatus = GL_FALSE;
}
}
extern "C" const unsigned *
brw_gs_emit(struct brw_context *brw,
struct gl_shader_program *prog,
struct brw_gs_compile *c,
void *mem_ctx,
unsigned *final_assembly_size)
{
if (unlikely(INTEL_DEBUG & DEBUG_GS)) {
struct brw_shader *shader =
(brw_shader *) prog->_LinkedShaders[MESA_SHADER_GEOMETRY];
brw_dump_ir(brw, "geometry", prog, &shader->base, NULL);
}
/* Compile the geometry shader in DUAL_OBJECT dispatch mode, if we can do
* so without spilling. If the GS invocations count > 1, then we can't use
* dual object mode.
*/
if (c->prog_data.invocations <= 1 &&
likely(!(INTEL_DEBUG & DEBUG_NO_DUAL_OBJECT_GS))) {
c->prog_data.dual_instanced_dispatch = false;
vec4_gs_visitor v(brw, c, prog, mem_ctx, true /* no_spills */);
if (v.run()) {
return generate_assembly(brw, prog, &c->gp->program.Base,
&c->prog_data.base, mem_ctx, v.cfg,
final_assembly_size);
}
}
/* Either we failed to compile in DUAL_OBJECT mode (probably because it
* would have required spilling) or DUAL_OBJECT mode is disabled. So fall
* back to DUAL_INSTANCED mode, which consumes fewer registers.
*
* FIXME: In an ideal world we'd fall back to SINGLE mode, which would
* allow us to interleave general purpose registers (resulting in even less
* likelihood of spilling). But at the moment, the vec4 generator and
* visitor classes don't have the infrastructure to interleave general
* purpose registers, so DUAL_INSTANCED is the best we can do.
*/
c->prog_data.dual_instanced_dispatch = true;
vec4_gs_visitor v(brw, c, prog, mem_ctx, false /* no_spills */);
if (!v.run()) {
prog->LinkStatus = false;
ralloc_strcat(&prog->InfoLog, v.fail_msg);
return NULL;
}
return generate_assembly(brw, prog, &c->gp->program.Base, &c->prog_data.base,
mem_ctx, v.cfg, final_assembly_size);
}
void
glcpp_warning (YYLTYPE *locp, glcpp_parser_t *parser, const char *fmt, ...)
{
va_list ap;
ralloc_asprintf_append(&parser->info_log, "%u:%u(%u): "
"preprocessor warning: ",
locp->source,
locp->first_line,
locp->first_column);
va_start(ap, fmt);
ralloc_vasprintf_append(&parser->info_log, fmt, ap);
va_end(ap);
ralloc_strcat(&parser->info_log, "\n");
}
/**
* For GL_EXT_separate_shader_objects
*/
GLuint GLAPIENTRY
_mesa_CreateShaderProgramEXT(GLenum type, const GLchar *string)
{
GET_CURRENT_CONTEXT(ctx);
const GLuint shader = create_shader(ctx, type);
GLuint program = 0;
if (shader) {
shader_source(ctx, shader, _mesa_strdup(string));
compile_shader(ctx, shader);
program = create_shader_program(ctx);
if (program) {
struct gl_shader_program *shProg;
struct gl_shader *sh;
GLint compiled = GL_FALSE;
shProg = _mesa_lookup_shader_program(ctx, program);
sh = _mesa_lookup_shader(ctx, shader);
get_shaderiv(ctx, shader, GL_COMPILE_STATUS, &compiled);
if (compiled) {
attach_shader(ctx, program, shader);
link_program(ctx, program);
detach_shader(ctx, program, shader);
#if 0
/* Possibly... */
if (active-user-defined-varyings-in-linked-program) {
append-error-to-info-log;
shProg->LinkStatus = GL_FALSE;
}
#endif
}
ralloc_strcat(&shProg->InfoLog, sh->InfoLog);
}
delete_shader(ctx, shader);
}
return program;
}
int
preprocess(void *ralloc_ctx, const char **shader, char **info_log)
{
int errors;
glcpp_parser_t *parser = glcpp_parser_create ();
*shader = remove_line_continuations(parser, *shader);
glcpp_lex_set_source_string (parser, *shader);
glcpp_parser_parse (parser);
if (parser->skip_stack)
glcpp_error (&parser->skip_stack->loc, parser, "Unterminated #if\n");
ralloc_strcat(info_log, parser->info_log);
*shader = ralloc_strdup(ralloc_ctx, parser->output);
errors = parser->error;
glcpp_parser_destroy (parser);
return errors;
}
/**
* All Mesa program -> GPU code generation goes through this function.
* Depending on the instructions used (i.e. flow control instructions)
* we'll use one of two code generators.
*/
bool
brw_codegen_wm_prog(struct brw_context *brw,
struct gl_shader_program *prog,
struct brw_fragment_program *fp,
struct brw_wm_prog_key *key)
{
struct gl_context *ctx = &brw->ctx;
void *mem_ctx = ralloc_context(NULL);
struct brw_wm_prog_data prog_data;
const GLuint *program;
struct brw_shader *fs = NULL;
GLuint program_size;
bool start_busy = false;
double start_time = 0;
if (prog)
fs = (struct brw_shader *)prog->_LinkedShaders[MESA_SHADER_FRAGMENT];
memset(&prog_data, 0, sizeof(prog_data));
/* Use ALT floating point mode for ARB programs so that 0^0 == 1. */
if (!prog)
prog_data.base.use_alt_mode = true;
assign_fs_binding_table_offsets(brw->intelScreen->devinfo, prog,
&fp->program.Base, key, &prog_data);
/* Allocate the references to the uniforms that will end up in the
* prog_data associated with the compiled program, and which will be freed
* by the state cache.
*/
int param_count = fp->program.Base.nir->num_uniforms;
if (fs)
prog_data.base.nr_image_params = fs->base.NumImages;
/* The backend also sometimes adds params for texture size. */
param_count += 2 * ctx->Const.Program[MESA_SHADER_FRAGMENT].MaxTextureImageUnits;
prog_data.base.param =
rzalloc_array(NULL, const gl_constant_value *, param_count);
prog_data.base.pull_param =
rzalloc_array(NULL, const gl_constant_value *, param_count);
prog_data.base.image_param =
rzalloc_array(NULL, struct brw_image_param,
prog_data.base.nr_image_params);
prog_data.base.nr_params = param_count;
if (prog) {
brw_nir_setup_glsl_uniforms(fp->program.Base.nir, prog, &fp->program.Base,
&prog_data.base, true);
} else {
brw_nir_setup_arb_uniforms(fp->program.Base.nir, &fp->program.Base,
&prog_data.base);
}
if (unlikely(brw->perf_debug)) {
start_busy = (brw->batch.last_bo &&
drm_intel_bo_busy(brw->batch.last_bo));
start_time = get_time();
}
if (unlikely(INTEL_DEBUG & DEBUG_WM))
brw_dump_ir("fragment", prog, fs ? &fs->base : NULL, &fp->program.Base);
int st_index8 = -1, st_index16 = -1;
if (INTEL_DEBUG & DEBUG_SHADER_TIME) {
st_index8 = brw_get_shader_time_index(brw, prog, &fp->program.Base, ST_FS8);
st_index16 = brw_get_shader_time_index(brw, prog, &fp->program.Base, ST_FS16);
}
char *error_str = NULL;
program = brw_compile_fs(brw->intelScreen->compiler, brw, mem_ctx,
key, &prog_data, fp->program.Base.nir,
&fp->program.Base, st_index8, st_index16,
brw->use_rep_send, &program_size, &error_str);
if (program == NULL) {
if (prog) {
prog->LinkStatus = false;
ralloc_strcat(&prog->InfoLog, error_str);
}
_mesa_problem(NULL, "Failed to compile fragment shader: %s\n", error_str);
ralloc_free(mem_ctx);
return false;
}
if (unlikely(brw->perf_debug) && fs) {
if (fs->compiled_once)
brw_wm_debug_recompile(brw, prog, key);
fs->compiled_once = true;
if (start_busy && !drm_intel_bo_busy(brw->batch.last_bo)) {
perf_debug("FS compile took %.03f ms and stalled the GPU\n",
(get_time() - start_time) * 1000);
}
}
if (prog_data.base.total_scratch) {
brw_get_scratch_bo(brw, &brw->wm.base.scratch_bo,
prog_data.base.total_scratch * brw->max_wm_threads);
}
if (unlikely(INTEL_DEBUG & DEBUG_WM))
fprintf(stderr, "\n");
brw_upload_cache(&brw->cache, BRW_CACHE_FS_PROG,
key, sizeof(struct brw_wm_prog_key),
program, program_size,
&prog_data, sizeof(prog_data),
&brw->wm.base.prog_offset, &brw->wm.prog_data);
ralloc_free(mem_ctx);
return true;
}
static const unsigned *
brw_cs_emit(struct brw_context *brw,
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,
unsigned *final_assembly_size)
{
bool start_busy = false;
double start_time = 0;
if (unlikely(brw->perf_debug)) {
start_busy = (brw->batch.last_bo &&
drm_intel_bo_busy(brw->batch.last_bo));
start_time = get_time();
}
struct brw_shader *shader =
(struct brw_shader *) prog->_LinkedShaders[MESA_SHADER_COMPUTE];
if (unlikely(INTEL_DEBUG & DEBUG_CS))
brw_dump_ir("compute", prog, &shader->base, &cp->Base);
prog_data->local_size[0] = cp->LocalSize[0];
prog_data->local_size[1] = cp->LocalSize[1];
prog_data->local_size[2] = cp->LocalSize[2];
unsigned local_workgroup_size =
cp->LocalSize[0] * cp->LocalSize[1] * cp->LocalSize[2];
cfg_t *cfg = NULL;
const char *fail_msg = NULL;
int st_index = -1;
if (INTEL_DEBUG & DEBUG_SHADER_TIME)
st_index = brw_get_shader_time_index(brw, prog, &cp->Base, ST_CS);
/* Now the main event: Visit the shader IR and generate our CS IR for it.
*/
fs_visitor v8(brw->intelScreen->compiler, brw,
mem_ctx, MESA_SHADER_COMPUTE, key, &prog_data->base, prog,
&cp->Base, 8, st_index);
if (!v8.run_cs()) {
fail_msg = v8.fail_msg;
} else if (local_workgroup_size <= 8 * brw->max_cs_threads) {
cfg = v8.cfg;
prog_data->simd_size = 8;
}
fs_visitor v16(brw->intelScreen->compiler, brw,
mem_ctx, MESA_SHADER_COMPUTE, key, &prog_data->base, prog,
&cp->Base, 16, st_index);
if (likely(!(INTEL_DEBUG & DEBUG_NO16)) &&
!fail_msg && !v8.simd16_unsupported &&
local_workgroup_size <= 16 * brw->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);
if (!cfg) {
fail_msg =
"Couldn't generate SIMD16 program and not "
"enough threads for SIMD8";
}
} else {
cfg = v16.cfg;
prog_data->simd_size = 16;
}
}
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);
return NULL;
}
fs_generator g(brw->intelScreen->compiler, brw,
mem_ctx, (void*) key, &prog_data->base, &cp->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);
g.enable_debug(name);
}
g.generate_code(cfg, prog_data->simd_size);
if (unlikely(brw->perf_debug) && shader) {
if (shader->compiled_once) {
_mesa_problem(&brw->ctx, "CS programs shouldn't need recompiles");
}
shader->compiled_once = true;
if (start_busy && !drm_intel_bo_busy(brw->batch.last_bo)) {
perf_debug("CS compile took %.03f ms and stalled the GPU\n",
(get_time() - start_time) * 1000);
}
}
return g.get_assembly(final_assembly_size);
}
static bool
brw_codegen_cs_prog(struct brw_context *brw,
struct gl_shader_program *prog,
struct brw_compute_program *cp,
struct brw_cs_prog_key *key)
{
struct gl_context *ctx = &brw->ctx;
const GLuint *program;
void *mem_ctx = ralloc_context(NULL);
GLuint program_size;
struct brw_cs_prog_data prog_data;
bool start_busy = false;
double start_time = 0;
struct brw_shader *cs =
(struct brw_shader *) prog->_LinkedShaders[MESA_SHADER_COMPUTE];
assert (cs);
memset(&prog_data, 0, sizeof(prog_data));
if (prog->Comp.SharedSize > 64 * 1024) {
prog->LinkStatus = false;
const char *error_str =
"Compute shader used more than 64KB of shared variables";
ralloc_strcat(&prog->InfoLog, error_str);
_mesa_problem(NULL, "Failed to link compute shader: %s\n", error_str);
ralloc_free(mem_ctx);
return false;
} else {
prog_data.base.total_shared = prog->Comp.SharedSize;
}
assign_cs_binding_table_offsets(brw->intelScreen->devinfo, prog,
&cp->program.Base, &prog_data);
/* Allocate the references to the uniforms that will end up in the
* prog_data associated with the compiled program, and which will be freed
* by the state cache.
*/
int param_count = cp->program.Base.nir->num_uniforms;
/* The backend also sometimes adds params for texture size. */
param_count += 2 * ctx->Const.Program[MESA_SHADER_COMPUTE].MaxTextureImageUnits;
prog_data.base.param =
rzalloc_array(NULL, const gl_constant_value *, param_count);
prog_data.base.pull_param =
rzalloc_array(NULL, const gl_constant_value *, param_count);
prog_data.base.image_param =
rzalloc_array(NULL, struct brw_image_param, cs->base.NumImages);
prog_data.base.nr_params = param_count;
prog_data.base.nr_image_params = cs->base.NumImages;
brw_nir_setup_glsl_uniforms(cp->program.Base.nir, prog, &cp->program.Base,
&prog_data.base, true);
if (unlikely(brw->perf_debug)) {
start_busy = (brw->batch.last_bo &&
drm_intel_bo_busy(brw->batch.last_bo));
start_time = get_time();
}
if (unlikely(INTEL_DEBUG & DEBUG_CS))
brw_dump_ir("compute", prog, &cs->base, &cp->program.Base);
int st_index = -1;
if (INTEL_DEBUG & DEBUG_SHADER_TIME)
st_index = brw_get_shader_time_index(brw, prog, &cp->program.Base, ST_CS);
char *error_str;
program = brw_compile_cs(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;
}
if (unlikely(brw->perf_debug) && cs) {
if (cs->compiled_once) {
_mesa_problem(&brw->ctx, "CS programs shouldn't need recompiles");
}
cs->compiled_once = true;
if (start_busy && !drm_intel_bo_busy(brw->batch.last_bo)) {
perf_debug("CS compile took %.03f ms and stalled the GPU\n",
(get_time() - start_time) * 1000);
}
}
if (prog_data.base.total_scratch) {
brw_get_scratch_bo(brw, &brw->cs.base.scratch_bo,
prog_data.base.total_scratch * brw->max_cs_threads);
}
if (unlikely(INTEL_DEBUG & DEBUG_CS))
fprintf(stderr, "\n");
brw_upload_cache(&brw->cache, BRW_CACHE_CS_PROG,
key, sizeof(*key),
program, program_size,
&prog_data, sizeof(prog_data),
&brw->cs.base.prog_offset, &brw->cs.prog_data);
ralloc_free(mem_ctx);
return true;
}
static bool
brw_codegen_gs_prog(struct brw_context *brw,
struct brw_program *gp,
struct brw_gs_prog_key *key)
{
struct brw_compiler *compiler = brw->screen->compiler;
const struct gen_device_info *devinfo = &brw->screen->devinfo;
struct brw_stage_state *stage_state = &brw->gs.base;
struct brw_gs_prog_data prog_data;
bool start_busy = false;
double start_time = 0;
memset(&prog_data, 0, sizeof(prog_data));
assign_gs_binding_table_offsets(devinfo, &gp->program, &prog_data);
/* Allocate the references to the uniforms that will end up in the
* prog_data associated with the compiled program, and which will be freed
* by the state cache.
*
* Note: param_count needs to be num_uniform_components * 4, since we add
* padding around uniform values below vec4 size, so the worst case is that
* every uniform is a float which gets padded to the size of a vec4.
*/
int param_count = gp->program.nir->num_uniforms / 4;
prog_data.base.base.param =
rzalloc_array(NULL, const gl_constant_value *, param_count);
prog_data.base.base.pull_param =
rzalloc_array(NULL, const gl_constant_value *, param_count);
prog_data.base.base.image_param =
rzalloc_array(NULL, struct brw_image_param,
gp->program.info.num_images);
prog_data.base.base.nr_params = param_count;
prog_data.base.base.nr_image_params = gp->program.info.num_images;
brw_nir_setup_glsl_uniforms(gp->program.nir, &gp->program,
&prog_data.base.base,
compiler->scalar_stage[MESA_SHADER_GEOMETRY]);
brw_nir_analyze_ubo_ranges(compiler, gp->program.nir,
prog_data.base.base.ubo_ranges);
uint64_t outputs_written = gp->program.info.outputs_written;
brw_compute_vue_map(devinfo,
&prog_data.base.vue_map, outputs_written,
gp->program.info.separate_shader);
int st_index = -1;
if (INTEL_DEBUG & DEBUG_SHADER_TIME)
st_index = brw_get_shader_time_index(brw, &gp->program, ST_GS, true);
if (unlikely(brw->perf_debug)) {
start_busy = brw->batch.last_bo && brw_bo_busy(brw->batch.last_bo);
start_time = get_time();
}
void *mem_ctx = ralloc_context(NULL);
unsigned program_size;
char *error_str;
const unsigned *program =
brw_compile_gs(brw->screen->compiler, brw, mem_ctx, key,
&prog_data, gp->program.nir, &gp->program,
st_index, &program_size, &error_str);
if (program == NULL) {
ralloc_strcat(&gp->program.sh.data->InfoLog, error_str);
_mesa_problem(NULL, "Failed to compile geometry shader: %s\n", error_str);
ralloc_free(mem_ctx);
return false;
}
if (unlikely(brw->perf_debug)) {
if (gp->compiled_once) {
brw_gs_debug_recompile(brw, &gp->program, key);
}
if (start_busy && !brw_bo_busy(brw->batch.last_bo)) {
perf_debug("GS compile took %.03f ms and stalled the GPU\n",
(get_time() - start_time) * 1000);
}
gp->compiled_once = true;
}
/* Scratch space is used for register spilling */
brw_alloc_stage_scratch(brw, stage_state,
prog_data.base.base.total_scratch,
devinfo->max_gs_threads);
brw_upload_cache(&brw->cache, BRW_CACHE_GS_PROG,
key, sizeof(*key),
program, program_size,
&prog_data, sizeof(prog_data),
&stage_state->prog_offset, &brw->gs.base.prog_data);
ralloc_free(mem_ctx);
return true;
}
static bool
brw_codegen_tes_prog(struct brw_context *brw,
struct brw_program *tep,
struct brw_tes_prog_key *key)
{
const struct brw_compiler *compiler = brw->screen->compiler;
const struct gen_device_info *devinfo = &brw->screen->devinfo;
struct brw_stage_state *stage_state = &brw->tes.base;
struct brw_tes_prog_data prog_data;
bool start_busy = false;
double start_time = 0;
memset(&prog_data, 0, sizeof(prog_data));
void *mem_ctx = ralloc_context(NULL);
nir_shader *nir = nir_shader_clone(mem_ctx, tep->program.nir);
brw_assign_common_binding_table_offsets(devinfo, &tep->program,
&prog_data.base.base, 0);
brw_nir_setup_glsl_uniforms(mem_ctx, nir, &tep->program,
&prog_data.base.base,
compiler->scalar_stage[MESA_SHADER_TESS_EVAL]);
brw_nir_analyze_ubo_ranges(compiler, nir, NULL,
prog_data.base.base.ubo_ranges);
int st_index = -1;
if (unlikely(INTEL_DEBUG & DEBUG_SHADER_TIME))
st_index = brw_get_shader_time_index(brw, &tep->program, ST_TES, true);
if (unlikely(brw->perf_debug)) {
start_busy = brw->batch.last_bo && brw_bo_busy(brw->batch.last_bo);
start_time = get_time();
}
struct brw_vue_map input_vue_map;
brw_compute_tess_vue_map(&input_vue_map, key->inputs_read,
key->patch_inputs_read);
char *error_str;
const unsigned *program =
brw_compile_tes(compiler, brw, mem_ctx, key, &input_vue_map, &prog_data,
nir, &tep->program, st_index, &error_str);
if (program == NULL) {
tep->program.sh.data->LinkStatus = LINKING_FAILURE;
ralloc_strcat(&tep->program.sh.data->InfoLog, error_str);
_mesa_problem(NULL, "Failed to compile tessellation evaluation shader: "
"%s\n", error_str);
ralloc_free(mem_ctx);
return false;
}
if (unlikely(brw->perf_debug)) {
if (tep->compiled_once) {
brw_debug_recompile(brw, MESA_SHADER_TESS_EVAL, tep->program.Id,
key->program_string_id, key);
}
if (start_busy && !brw_bo_busy(brw->batch.last_bo)) {
perf_debug("TES compile took %.03f ms and stalled the GPU\n",
(get_time() - start_time) * 1000);
}
tep->compiled_once = true;
}
/* Scratch space is used for register spilling */
brw_alloc_stage_scratch(brw, stage_state,
prog_data.base.base.total_scratch);
/* The param and pull_param arrays will be freed by the shader cache. */
ralloc_steal(NULL, prog_data.base.base.param);
ralloc_steal(NULL, prog_data.base.base.pull_param);
brw_upload_cache(&brw->cache, BRW_CACHE_TES_PROG,
key, sizeof(*key),
program, prog_data.base.base.program_size,
&prog_data, sizeof(prog_data),
&stage_state->prog_offset, &brw->tes.base.prog_data);
ralloc_free(mem_ctx);
return true;
}
static bool
brw_codegen_cs_prog(struct brw_context *brw,
struct gl_shader_program *prog,
struct brw_program *cp,
struct brw_cs_prog_key *key)
{
const struct gen_device_info *devinfo = &brw->screen->devinfo;
struct gl_context *ctx = &brw->ctx;
const GLuint *program;
void *mem_ctx = ralloc_context(NULL);
GLuint program_size;
struct brw_cs_prog_data prog_data;
bool start_busy = false;
double start_time = 0;
struct brw_shader *cs =
(struct brw_shader *) prog->_LinkedShaders[MESA_SHADER_COMPUTE];
assert (cs);
memset(&prog_data, 0, sizeof(prog_data));
if (prog->Comp.SharedSize > 64 * 1024) {
prog->LinkStatus = false;
const char *error_str =
"Compute shader used more than 64KB of shared variables";
ralloc_strcat(&prog->InfoLog, error_str);
_mesa_problem(NULL, "Failed to link compute shader: %s\n", error_str);
ralloc_free(mem_ctx);
return false;
} else {
prog_data.base.total_shared = prog->Comp.SharedSize;
}
assign_cs_binding_table_offsets(devinfo, prog, &cp->program, &prog_data);
/* Allocate the references to the uniforms that will end up in the
* prog_data associated with the compiled program, and which will be freed
* by the state cache.
*/
int param_count = cp->program.nir->num_uniforms / 4;
/* The backend also sometimes add a param for the thread local id. */
prog_data.thread_local_id_index = param_count++;
/* The backend also sometimes adds params for texture size. */
param_count += 2 * ctx->Const.Program[MESA_SHADER_COMPUTE].MaxTextureImageUnits;
prog_data.base.param =
rzalloc_array(NULL, const gl_constant_value *, param_count);
prog_data.base.pull_param =
rzalloc_array(NULL, const gl_constant_value *, param_count);
prog_data.base.image_param =
rzalloc_array(NULL, struct brw_image_param, cs->base.NumImages);
prog_data.base.nr_params = param_count;
prog_data.base.nr_image_params = cs->base.NumImages;
brw_nir_setup_glsl_uniforms(cp->program.nir, prog, &cp->program,
&prog_data.base, true);
if (unlikely(brw->perf_debug)) {
start_busy = (brw->batch.last_bo &&
drm_intel_bo_busy(brw->batch.last_bo));
start_time = get_time();
}
if (unlikely(INTEL_DEBUG & DEBUG_CS))
brw_dump_ir("compute", prog, &cs->base, &cp->program);
int st_index = -1;
if (INTEL_DEBUG & DEBUG_SHADER_TIME)
st_index = brw_get_shader_time_index(brw, prog, &cp->program, ST_CS);
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);
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;
}
if (unlikely(brw->perf_debug) && cs) {
if (cs->compiled_once) {
_mesa_problem(&brw->ctx, "CS programs shouldn't need recompiles");
}
cs->compiled_once = true;
if (start_busy && !drm_intel_bo_busy(brw->batch.last_bo)) {
perf_debug("CS compile took %.03f ms and stalled the GPU\n",
(get_time() - start_time) * 1000);
}
}
const unsigned subslices = MAX2(brw->screen->subslice_total, 1);
/* WaCSScratchSize:hsw
*
* Haswell's scratch space address calculation appears to be sparse
* rather than tightly packed. The Thread ID has bits indicating
* which subslice, EU within a subslice, and thread within an EU
* it is. There's a maximum of two slices and two subslices, so these
* can be stored with a single bit. Even though there are only 10 EUs
* per subslice, this is stored in 4 bits, so there's an effective
* maximum value of 16 EUs. Similarly, although there are only 7
* threads per EU, this is stored in a 3 bit number, giving an effective
* maximum value of 8 threads per EU.
*
* This means that we need to use 16 * 8 instead of 10 * 7 for the
* number of threads per subslice.
*/
const unsigned scratch_ids_per_subslice =
brw->is_haswell ? 16 * 8 : devinfo->max_cs_threads;
brw_alloc_stage_scratch(brw, &brw->cs.base,
prog_data.base.total_scratch,
scratch_ids_per_subslice * subslices);
if (unlikely(INTEL_DEBUG & DEBUG_CS))
fprintf(stderr, "\n");
brw_upload_cache(&brw->cache, BRW_CACHE_CS_PROG,
key, sizeof(*key),
program, program_size,
&prog_data, sizeof(prog_data),
&brw->cs.base.prog_offset, &brw->cs.base.prog_data);
ralloc_free(mem_ctx);
return true;
}
bool
shader_cache_read_program_metadata(struct gl_context *ctx,
struct gl_shader_program *prog)
{
/* Fixed function programs generated by Mesa are not cached. So don't
* try to read metadata for them from the cache.
*/
if (prog->Name == 0)
return false;
struct disk_cache *cache = ctx->Cache;
if (!cache)
return false;
/* Include bindings when creating sha1. These bindings change the resulting
* binary so they are just as important as the shader source.
*/
char *buf = ralloc_strdup(NULL, "vb: ");
prog->AttributeBindings->iterate(create_binding_str, &buf);
ralloc_strcat(&buf, "fb: ");
prog->FragDataBindings->iterate(create_binding_str, &buf);
ralloc_strcat(&buf, "fbi: ");
prog->FragDataIndexBindings->iterate(create_binding_str, &buf);
/* SSO has an effect on the linked program so include this when generating
* the sha also.
*/
ralloc_asprintf_append(&buf, "sso: %s\n",
prog->SeparateShader ? "T" : "F");
/* A shader might end up producing different output depending on the glsl
* version supported by the compiler. For example a different path might be
* taken by the preprocessor, so add the version to the hash input.
*/
ralloc_asprintf_append(&buf, "api: %d glsl: %d fglsl: %d\n",
ctx->API, ctx->Const.GLSLVersion,
ctx->Const.ForceGLSLVersion);
/* We run the preprocessor on shaders after hashing them, so we need to
* add any extension override vars to the hash. If we don't do this the
* preprocessor could result in different output and we could load the
* wrong shader.
*/
char *ext_override = getenv("MESA_EXTENSION_OVERRIDE");
if (ext_override) {
ralloc_asprintf_append(&buf, "ext:%s", ext_override);
}
/* DRI config options may also change the output from the compiler so
* include them as an input to sha1 creation.
*/
char sha1buf[41];
_mesa_sha1_format(sha1buf, ctx->Const.dri_config_options_sha1);
ralloc_strcat(&buf, sha1buf);
for (unsigned i = 0; i < prog->NumShaders; i++) {
struct gl_shader *sh = prog->Shaders[i];
_mesa_sha1_format(sha1buf, sh->sha1);
ralloc_asprintf_append(&buf, "%s: %s\n",
_mesa_shader_stage_to_abbrev(sh->Stage), sha1buf);
}
disk_cache_compute_key(cache, buf, strlen(buf), prog->data->sha1);
ralloc_free(buf);
size_t size;
uint8_t *buffer = (uint8_t *) disk_cache_get(cache, prog->data->sha1,
&size);
if (buffer == NULL) {
/* Cached program not found. We may have seen the individual shaders
* before and skipped compiling but they may not have been used together
* in this combination before. Fall back to linking shaders but first
* re-compile the shaders.
*
* We could probably only compile the shaders which were skipped here
* but we need to be careful because the source may also have been
* changed since the last compile so for now we just recompile
* everything.
*/
compile_shaders(ctx, prog);
return false;
}
if (ctx->_Shader->Flags & GLSL_CACHE_INFO) {
_mesa_sha1_format(sha1buf, prog->data->sha1);
fprintf(stderr, "loading shader program meta data from cache: %s\n",
sha1buf);
}
struct blob_reader metadata;
blob_reader_init(&metadata, buffer, size);
bool deserialized = deserialize_glsl_program(&metadata, ctx, prog);
if (!deserialized || metadata.current != metadata.end || metadata.overrun) {
/* Something has gone wrong discard the item from the cache and rebuild
* from source.
*/
assert(!"Invalid GLSL shader disk cache item!");
if (ctx->_Shader->Flags & GLSL_CACHE_INFO) {
fprintf(stderr, "Error reading program from cache (invalid GLSL "
"cache item)\n");
}
disk_cache_remove(cache, prog->data->sha1);
compile_shaders(ctx, prog);
free(buffer);
return false;
}
/* This is used to flag a shader retrieved from cache */
prog->data->LinkStatus = linking_skipped;
/* Since the program load was successful, CompileStatus of all shaders at
* this point should normally be compile_skipped. However because of how
* the eviction works, it may happen that some of the individual shader keys
* have been evicted, resulting in unnecessary recompiles on this load, so
* mark them again to skip such recompiles next time.
*/
char sha1_buf[41];
for (unsigned i = 0; i < prog->NumShaders; i++) {
if (prog->Shaders[i]->CompileStatus == compiled_no_opts) {
disk_cache_put_key(cache, prog->Shaders[i]->sha1);
if (ctx->_Shader->Flags & GLSL_CACHE_INFO) {
_mesa_sha1_format(sha1_buf, prog->Shaders[i]->sha1);
fprintf(stderr, "re-marking shader: %s\n", sha1_buf);
}
}
}
free (buffer);
return true;
}
bool
brw_codegen_vs_prog(struct brw_context *brw,
struct gl_shader_program *prog,
struct brw_vertex_program *vp,
struct brw_vs_prog_key *key)
{
GLuint program_size;
const GLuint *program;
struct brw_vs_prog_data prog_data;
struct brw_stage_prog_data *stage_prog_data = &prog_data.base.base;
void *mem_ctx;
int i;
struct brw_shader *vs = NULL;
bool start_busy = false;
double start_time = 0;
if (prog)
vs = (struct brw_shader *) prog->_LinkedShaders[MESA_SHADER_VERTEX];
memset(&prog_data, 0, sizeof(prog_data));
/* Use ALT floating point mode for ARB programs so that 0^0 == 1. */
if (!prog)
stage_prog_data->use_alt_mode = true;
mem_ctx = ralloc_context(NULL);
brw_assign_common_binding_table_offsets(MESA_SHADER_VERTEX,
brw->intelScreen->devinfo,
prog, &vp->program.Base,
&prog_data.base.base, 0);
/* Allocate the references to the uniforms that will end up in the
* prog_data associated with the compiled program, and which will be freed
* by the state cache.
*/
int param_count = vp->program.Base.nir->num_uniforms;
if (!brw->intelScreen->compiler->scalar_vs)
param_count *= 4;
if (vs)
prog_data.base.base.nr_image_params = vs->base.NumImages;
/* vec4_visitor::setup_uniform_clipplane_values() also uploads user clip
* planes as uniforms.
*/
param_count += key->nr_userclip_plane_consts * 4;
stage_prog_data->param =
rzalloc_array(NULL, const gl_constant_value *, param_count);
stage_prog_data->pull_param =
rzalloc_array(NULL, const gl_constant_value *, param_count);
stage_prog_data->image_param =
rzalloc_array(NULL, struct brw_image_param,
stage_prog_data->nr_image_params);
stage_prog_data->nr_params = param_count;
if (prog) {
brw_nir_setup_glsl_uniforms(vp->program.Base.nir, prog, &vp->program.Base,
&prog_data.base.base,
brw->intelScreen->compiler->scalar_vs);
} else {
brw_nir_setup_arb_uniforms(vp->program.Base.nir, &vp->program.Base,
&prog_data.base.base);
}
GLbitfield64 outputs_written = vp->program.Base.OutputsWritten;
prog_data.inputs_read = vp->program.Base.InputsRead;
if (key->copy_edgeflag) {
outputs_written |= BITFIELD64_BIT(VARYING_SLOT_EDGE);
prog_data.inputs_read |= VERT_BIT_EDGEFLAG;
}
if (brw->gen < 6) {
/* Put dummy slots into the VUE for the SF to put the replaced
* point sprite coords in. We shouldn't need these dummy slots,
* which take up precious URB space, but it would mean that the SF
* doesn't get nice aligned pairs of input coords into output
* coords, which would be a pain to handle.
*/
for (i = 0; i < 8; i++) {
if (key->point_coord_replace & (1 << i))
outputs_written |= BITFIELD64_BIT(VARYING_SLOT_TEX0 + i);
}
/* if back colors are written, allocate slots for front colors too */
if (outputs_written & BITFIELD64_BIT(VARYING_SLOT_BFC0))
outputs_written |= BITFIELD64_BIT(VARYING_SLOT_COL0);
if (outputs_written & BITFIELD64_BIT(VARYING_SLOT_BFC1))
outputs_written |= BITFIELD64_BIT(VARYING_SLOT_COL1);
}
/* In order for legacy clipping to work, we need to populate the clip
* distance varying slots whenever clipping is enabled, even if the vertex
* shader doesn't write to gl_ClipDistance.
*/
if (key->nr_userclip_plane_consts > 0) {
outputs_written |= BITFIELD64_BIT(VARYING_SLOT_CLIP_DIST0);
outputs_written |= BITFIELD64_BIT(VARYING_SLOT_CLIP_DIST1);
}
brw_compute_vue_map(brw->intelScreen->devinfo,
&prog_data.base.vue_map, outputs_written,
prog ? prog->SeparateShader : false);
if (0) {
_mesa_fprint_program_opt(stderr, &vp->program.Base, PROG_PRINT_DEBUG,
true);
}
if (unlikely(brw->perf_debug)) {
start_busy = (brw->batch.last_bo &&
drm_intel_bo_busy(brw->batch.last_bo));
start_time = get_time();
}
if (unlikely(INTEL_DEBUG & DEBUG_VS))
brw_dump_ir("vertex", prog, vs ? &vs->base : NULL, &vp->program.Base);
int st_index = -1;
if (INTEL_DEBUG & DEBUG_SHADER_TIME)
st_index = brw_get_shader_time_index(brw, prog, &vp->program.Base, ST_VS);
/* Emit GEN4 code.
*/
char *error_str;
program = brw_compile_vs(brw->intelScreen->compiler, brw, mem_ctx, key,
&prog_data, vp->program.Base.nir,
brw_select_clip_planes(&brw->ctx),
!_mesa_is_gles3(&brw->ctx),
st_index, &program_size, &error_str);
if (program == NULL) {
if (prog) {
prog->LinkStatus = false;
ralloc_strcat(&prog->InfoLog, error_str);
}
_mesa_problem(NULL, "Failed to compile vertex shader: %s\n", error_str);
ralloc_free(mem_ctx);
return false;
}
if (unlikely(brw->perf_debug) && vs) {
if (vs->compiled_once) {
brw_vs_debug_recompile(brw, prog, key);
}
if (start_busy && !drm_intel_bo_busy(brw->batch.last_bo)) {
perf_debug("VS compile took %.03f ms and stalled the GPU\n",
(get_time() - start_time) * 1000);
}
vs->compiled_once = true;
}
/* Scratch space is used for register spilling */
if (prog_data.base.base.total_scratch) {
brw_get_scratch_bo(brw, &brw->vs.base.scratch_bo,
prog_data.base.base.total_scratch *
brw->max_vs_threads);
}
brw_upload_cache(&brw->cache, BRW_CACHE_VS_PROG,
key, sizeof(struct brw_vs_prog_key),
program, program_size,
&prog_data, sizeof(prog_data),
&brw->vs.base.prog_offset, &brw->vs.prog_data);
ralloc_free(mem_ctx);
return true;
}
static bool
brw_codegen_wm_prog(struct brw_context *brw,
struct brw_program *fp,
struct brw_wm_prog_key *key,
struct brw_vue_map *vue_map)
{
const struct gen_device_info *devinfo = &brw->screen->devinfo;
void *mem_ctx = ralloc_context(NULL);
struct brw_wm_prog_data prog_data;
const GLuint *program;
bool start_busy = false;
double start_time = 0;
nir_shader *nir = nir_shader_clone(mem_ctx, fp->program.nir);
memset(&prog_data, 0, sizeof(prog_data));
/* Use ALT floating point mode for ARB programs so that 0^0 == 1. */
if (fp->program.is_arb_asm)
prog_data.base.use_alt_mode = true;
assign_fs_binding_table_offsets(devinfo, &fp->program, key, &prog_data);
if (!fp->program.is_arb_asm) {
brw_nir_setup_glsl_uniforms(mem_ctx, nir, &fp->program,
&prog_data.base, true);
brw_nir_analyze_ubo_ranges(brw->screen->compiler, nir,
NULL, prog_data.base.ubo_ranges);
} else {
brw_nir_setup_arb_uniforms(mem_ctx, nir, &fp->program, &prog_data.base);
if (unlikely(INTEL_DEBUG & DEBUG_WM))
brw_dump_arb_asm("fragment", &fp->program);
}
if (unlikely(brw->perf_debug)) {
start_busy = (brw->batch.last_bo &&
brw_bo_busy(brw->batch.last_bo));
start_time = get_time();
}
int st_index8 = -1, st_index16 = -1, st_index32 = -1;
if (INTEL_DEBUG & DEBUG_SHADER_TIME) {
st_index8 = brw_get_shader_time_index(brw, &fp->program, ST_FS8,
!fp->program.is_arb_asm);
st_index16 = brw_get_shader_time_index(brw, &fp->program, ST_FS16,
!fp->program.is_arb_asm);
st_index32 = brw_get_shader_time_index(brw, &fp->program, ST_FS32,
!fp->program.is_arb_asm);
}
char *error_str = NULL;
program = brw_compile_fs(brw->screen->compiler, brw, mem_ctx,
key, &prog_data, nir,
&fp->program, st_index8, st_index16, st_index32,
true, false, vue_map,
&error_str);
if (program == NULL) {
if (!fp->program.is_arb_asm) {
fp->program.sh.data->LinkStatus = LINKING_FAILURE;
ralloc_strcat(&fp->program.sh.data->InfoLog, error_str);
}
_mesa_problem(NULL, "Failed to compile fragment shader: %s\n", error_str);
ralloc_free(mem_ctx);
return false;
}
if (unlikely(brw->perf_debug)) {
if (fp->compiled_once) {
brw_debug_recompile(brw, MESA_SHADER_FRAGMENT, fp->program.Id,
key->program_string_id, key);
}
fp->compiled_once = true;
if (start_busy && !brw_bo_busy(brw->batch.last_bo)) {
perf_debug("FS compile took %.03f ms and stalled the GPU\n",
(get_time() - start_time) * 1000);
}
}
brw_alloc_stage_scratch(brw, &brw->wm.base, prog_data.base.total_scratch);
if (unlikely((INTEL_DEBUG & DEBUG_WM) && fp->program.is_arb_asm))
fprintf(stderr, "\n");
/* The param and pull_param arrays will be freed by the shader cache. */
ralloc_steal(NULL, prog_data.base.param);
ralloc_steal(NULL, prog_data.base.pull_param);
brw_upload_cache(&brw->cache, BRW_CACHE_FS_PROG,
key, sizeof(struct brw_wm_prog_key),
program, prog_data.base.program_size,
&prog_data, sizeof(prog_data),
&brw->wm.base.prog_offset, &brw->wm.base.prog_data);
ralloc_free(mem_ctx);
return true;
}
static void
lower_sampler(nir_tex_instr *instr, struct gl_shader_program *shader_program,
const struct gl_program *prog, void *mem_ctx)
{
if (instr->sampler == NULL)
return;
/* Get the name and the offset */
instr->sampler_index = 0;
bool has_indirect = false;
char *name = ralloc_strdup(mem_ctx, instr->sampler->var->name);
for (nir_deref *deref = &instr->sampler->deref;
deref->child; deref = deref->child) {
switch (deref->child->deref_type) {
case nir_deref_type_array: {
nir_deref_array *deref_array = nir_deref_as_array(deref->child);
/* XXX: We're assuming here that the indirect is the last array
* thing we have. This should be ok for now as we don't support
* arrays_of_arrays yet.
*/
assert(!has_indirect);
instr->sampler_index *= glsl_get_length(deref->type);
switch (deref_array->deref_array_type) {
case nir_deref_array_type_direct:
instr->sampler_index += deref_array->base_offset;
if (deref_array->deref.child)
ralloc_asprintf_append(&name, "[%u]", deref_array->base_offset);
break;
case nir_deref_array_type_indirect: {
assert(!has_indirect);
instr->src = reralloc(mem_ctx, instr->src, nir_tex_src,
instr->num_srcs + 1);
memset(&instr->src[instr->num_srcs], 0, sizeof *instr->src);
instr->src[instr->num_srcs].src_type = nir_tex_src_sampler_offset;
instr->num_srcs++;
nir_instr_rewrite_src(&instr->instr,
&instr->src[instr->num_srcs - 1].src,
deref_array->indirect);
instr->sampler_array_size = glsl_get_length(deref->type);
nir_src empty;
memset(&empty, 0, sizeof empty);
nir_instr_rewrite_src(&instr->instr, &deref_array->indirect, empty);
if (deref_array->deref.child)
ralloc_strcat(&name, "[0]");
break;
}
case nir_deref_array_type_wildcard:
unreachable("Cannot copy samplers");
default:
unreachable("Invalid deref array type");
}
break;
}
case nir_deref_type_struct: {
nir_deref_struct *deref_struct = nir_deref_as_struct(deref->child);
const char *field = glsl_get_struct_elem_name(deref->type,
deref_struct->index);
ralloc_asprintf_append(&name, ".%s", field);
break;
}
default:
unreachable("Invalid deref type");
break;
}
}
instr->sampler_index += get_sampler_index(shader_program, name, prog);
instr->sampler = NULL;
}
