failed = true;
msg = ralloc_vasprintf(mem_ctx, format, va);
- msg = ralloc_asprintf(mem_ctx, "%s compile failed: %s\n", stage_abbrev, msg);
+ msg = ralloc_asprintf(mem_ctx, "SIMD%d %s compile failed: %s\n",
+ dispatch_width, stage_abbrev, msg);
this->fail_msg = msg;
break;
}
- if (inst->src[0].file == IMM) {
- assert(inst->src[0].type == BRW_REGISTER_TYPE_F);
+ break;
+ case BRW_OPCODE_ADD:
+ if (inst->src[1].file != IMM)
+ continue;
+
+ if (brw_reg_type_is_integer(inst->src[1].type) &&
+ inst->src[1].is_zero()) {
inst->opcode = BRW_OPCODE_MOV;
- inst->src[0].f *= inst->src[1].f;
inst->src[1] = reg_undef;
progress = true;
break;
}
- break;
- case BRW_OPCODE_ADD:
- if (inst->src[1].file != IMM)
- continue;
if (inst->src[0].file == IMM) {
assert(inst->src[0].type == BRW_REGISTER_TYPE_F);
return progress;
}
-/**
- * Optimize sample messages which are followed by the final RT write.
- *
- * CHV, and GEN9+ can mark a texturing SEND instruction with EOT to have its
- * results sent directly to the framebuffer, bypassing the EU. Recognize the
- * final texturing results copied to the framebuffer write payload and modify
- * them to write to the framebuffer directly.
- */
-bool
-fs_visitor::opt_sampler_eot()
-{
- brw_wm_prog_key *key = (brw_wm_prog_key*) this->key;
-
- if (stage != MESA_SHADER_FRAGMENT || dispatch_width > 16)
- return false;
-
- if (devinfo->gen != 9 && !devinfo->is_cherryview)
- return false;
-
- /* FINISHME: It should be possible to implement this optimization when there
- * are multiple drawbuffers.
- */
- if (key->nr_color_regions != 1)
- return false;
-
- /* Requires emitting a bunch of saturating MOV instructions during logical
- * send lowering to clamp the color payload, which the sampler unit isn't
- * going to do for us.
- */
- if (key->clamp_fragment_color)
- return false;
-
- /* Look for a texturing instruction immediately before the final FB_WRITE. */
- bblock_t *block = cfg->blocks[cfg->num_blocks - 1];
- fs_inst *fb_write = (fs_inst *)block->end();
- assert(fb_write->eot);
- assert(fb_write->opcode == FS_OPCODE_FB_WRITE_LOGICAL);
-
- /* There wasn't one; nothing to do. */
- if (unlikely(fb_write->prev->is_head_sentinel()))
- return false;
-
- fs_inst *tex_inst = (fs_inst *) fb_write->prev;
-
- /* 3D Sampler » Messages » Message Format
- *
- * “Response Length of zero is allowed on all SIMD8* and SIMD16* sampler
- * messages except sample+killpix, resinfo, sampleinfo, LOD, and gather4*”
- */
- if (tex_inst->opcode != SHADER_OPCODE_TEX_LOGICAL &&
- tex_inst->opcode != SHADER_OPCODE_TXD_LOGICAL &&
- tex_inst->opcode != SHADER_OPCODE_TXF_LOGICAL &&
- tex_inst->opcode != SHADER_OPCODE_TXL_LOGICAL &&
- tex_inst->opcode != FS_OPCODE_TXB_LOGICAL &&
- tex_inst->opcode != SHADER_OPCODE_TXF_CMS_LOGICAL &&
- tex_inst->opcode != SHADER_OPCODE_TXF_CMS_W_LOGICAL &&
- tex_inst->opcode != SHADER_OPCODE_TXF_UMS_LOGICAL)
- return false;
-
- /* XXX - This shouldn't be necessary. */
- if (tex_inst->prev->is_head_sentinel())
- return false;
-
- /* Check that the FB write sources are fully initialized by the single
- * texturing instruction.
- */
- for (unsigned i = 0; i < FB_WRITE_LOGICAL_NUM_SRCS; i++) {
- if (i == FB_WRITE_LOGICAL_SRC_COLOR0) {
- if (!fb_write->src[i].equals(tex_inst->dst) ||
- fb_write->size_read(i) != tex_inst->size_written)
- return false;
- } else if (i != FB_WRITE_LOGICAL_SRC_COMPONENTS) {
- if (fb_write->src[i].file != BAD_FILE)
- return false;
- }
- }
-
- assert(!tex_inst->eot); /* We can't get here twice */
- assert((tex_inst->offset & (0xff << 24)) == 0);
-
- const fs_builder ibld(this, block, tex_inst);
-
- tex_inst->offset |= fb_write->target << 24;
- tex_inst->eot = true;
- tex_inst->dst = ibld.null_reg_ud();
- tex_inst->size_written = 0;
- fb_write->remove(cfg->blocks[cfg->num_blocks - 1]);
-
- /* Marking EOT is sufficient, lower_logical_sends() will notice the EOT
- * flag and submit a header together with the sampler message as required
- * by the hardware.
- */
- invalidate_analysis(DEPENDENCY_INSTRUCTIONS | DEPENDENCY_VARIABLES);
- return true;
-}
-
bool
fs_visitor::opt_register_renaming()
{
high.offset = inst->dst.offset % REG_SIZE;
if (devinfo->gen >= 7) {
- if (inst->src[1].abs)
+ /* From GEN:BUG:1604601757:
+ *
+ * "When multiplying a DW and any lower precision integer, source modifier
+ * is not supported."
+ *
+ * An unsupported negate modifier on src[1] would ordinarily be
+ * lowered by the subsequent lower_regioning pass. In this case that
+ * pass would spawn another dword multiply. Instead, lower the
+ * modifier first.
+ */
+ const bool source_mods_unsupported = (devinfo->gen >= 12);
+
+ if (inst->src[1].abs || (inst->src[1].negate &&
+ source_mods_unsupported))
lower_src_modifiers(this, block, inst, 1);
if (inst->src[1].file == IMM) {
case FS_OPCODE_PACK:
case SHADER_OPCODE_SEL_EXEC:
case SHADER_OPCODE_CLUSTER_BROADCAST:
+ case SHADER_OPCODE_MOV_RELOC_IMM:
return get_fpu_lowered_simd_width(devinfo, inst);
case BRW_OPCODE_CMP: {
assert(dispatch_width % payload_width == 0);
assert(devinfo->gen >= 6);
- prog_data->uses_src_depth = prog_data->uses_src_w =
- (nir->info.system_values_read & (1ull << SYSTEM_VALUE_FRAG_COORD)) != 0;
-
- prog_data->uses_sample_mask =
- (nir->info.system_values_read & SYSTEM_BIT_SAMPLE_MASK_IN) != 0;
-
- /* From the Ivy Bridge PRM documentation for 3DSTATE_PS:
- *
- * "MSDISPMODE_PERSAMPLE is required in order to select
- * POSOFFSET_SAMPLE"
- *
- * So we can only really get sample positions if we are doing real
- * per-sample dispatch. If we need gl_SamplePosition and we don't have
- * persample dispatch, we hard-code it to 0.5.
- */
- prog_data->uses_pos_offset = prog_data->persample_dispatch &&
- (nir->info.system_values_read & SYSTEM_BIT_SAMPLE_POS);
-
/* R0: PS thread payload header. */
payload.num_regs++;
OPT(lower_simd_width);
OPT(lower_barycentrics);
-
- /* After SIMD lowering just in case we had to unroll the EOT send. */
- OPT(opt_sampler_eot);
-
OPT(lower_logical_sends);
/* After logical SEND lowering. */
/* Scheduling may create additional opportunities for CMOD propagation,
* so let's do it again. If CMOD propagation made any progress,
- * elminate dead code one more time.
+ * eliminate dead code one more time.
*/
bool progress = false;
const int iteration = 99;
prog_data->total_scratch = brw_get_scratch_size(last_scratch);
- if (stage == MESA_SHADER_COMPUTE) {
+ if (stage == MESA_SHADER_COMPUTE || stage == MESA_SHADER_KERNEL) {
if (devinfo->is_haswell) {
/* According to the MEDIA_VFE_STATE's "Per Thread Scratch Space"
* field documentation, Haswell supports a minimum of 2kB of
bool
fs_visitor::run_cs(bool allow_spilling)
{
- assert(stage == MESA_SHADER_COMPUTE);
+ assert(stage == MESA_SHADER_COMPUTE || stage == MESA_SHADER_KERNEL);
setup_cs_payload();
{
prog_data->flat_inputs = 0;
- nir_foreach_variable(var, &shader->inputs) {
+ nir_foreach_shader_in_variable(var, shader) {
unsigned slots = glsl_count_attribute_slots(var->type, false);
for (unsigned s = 0; s < slots; s++) {
int input_index = prog_data->urb_setup[var->data.location + s];
*
* This should be replaced by global value numbering someday.
*/
-static bool
-move_interpolation_to_top(nir_shader *nir)
+bool
+brw_nir_move_interpolation_to_top(nir_shader *nir)
{
bool progress = false;
}
}
}
- nir_metadata_preserve(f->impl, (nir_metadata)
- ((unsigned) nir_metadata_block_index |
- (unsigned) nir_metadata_dominance));
+ nir_metadata_preserve(f->impl, nir_metadata_block_index |
+ nir_metadata_dominance);
}
return progress;
*
* Useful when rendering to a non-multisampled buffer.
*/
-static bool
-demote_sample_qualifiers(nir_shader *nir)
+bool
+brw_nir_demote_sample_qualifiers(nir_shader *nir)
{
bool progress = true;
}
}
- nir_metadata_preserve(f->impl, (nir_metadata)
- ((unsigned) nir_metadata_block_index |
- (unsigned) nir_metadata_dominance));
+ nir_metadata_preserve(f->impl, nir_metadata_block_index |
+ nir_metadata_dominance);
}
return progress;
}
+void
+brw_nir_populate_wm_prog_data(const nir_shader *shader,
+ const struct gen_device_info *devinfo,
+ const struct brw_wm_prog_key *key,
+ struct brw_wm_prog_data *prog_data)
+{
+ prog_data->uses_src_depth = prog_data->uses_src_w =
+ shader->info.system_values_read & BITFIELD64_BIT(SYSTEM_VALUE_FRAG_COORD);
+
+ /* key->alpha_test_func means simulating alpha testing via discards,
+ * so the shader definitely kills pixels.
+ */
+ prog_data->uses_kill = shader->info.fs.uses_discard ||
+ key->alpha_test_func;
+ prog_data->uses_omask = !key->ignore_sample_mask_out &&
+ (shader->info.outputs_written & BITFIELD64_BIT(FRAG_RESULT_SAMPLE_MASK));
+ prog_data->computed_depth_mode = computed_depth_mode(shader);
+ prog_data->computed_stencil =
+ shader->info.outputs_written & BITFIELD64_BIT(FRAG_RESULT_STENCIL);
+
+ prog_data->persample_dispatch =
+ key->multisample_fbo &&
+ (key->persample_interp ||
+ (shader->info.system_values_read & (SYSTEM_BIT_SAMPLE_ID |
+ SYSTEM_BIT_SAMPLE_POS)) ||
+ shader->info.fs.uses_sample_qualifier ||
+ shader->info.outputs_read);
+
+ if (devinfo->gen >= 6) {
+ prog_data->uses_sample_mask =
+ shader->info.system_values_read & SYSTEM_BIT_SAMPLE_MASK_IN;
+
+ /* From the Ivy Bridge PRM documentation for 3DSTATE_PS:
+ *
+ * "MSDISPMODE_PERSAMPLE is required in order to select
+ * POSOFFSET_SAMPLE"
+ *
+ * So we can only really get sample positions if we are doing real
+ * per-sample dispatch. If we need gl_SamplePosition and we don't have
+ * persample dispatch, we hard-code it to 0.5.
+ */
+ prog_data->uses_pos_offset = prog_data->persample_dispatch &&
+ (shader->info.system_values_read & SYSTEM_BIT_SAMPLE_POS);
+ }
+
+ prog_data->has_render_target_reads = shader->info.outputs_read != 0ull;
+
+ prog_data->early_fragment_tests = shader->info.fs.early_fragment_tests;
+ prog_data->post_depth_coverage = shader->info.fs.post_depth_coverage;
+ prog_data->inner_coverage = shader->info.fs.inner_coverage;
+
+ prog_data->barycentric_interp_modes =
+ brw_compute_barycentric_interp_modes(devinfo, shader);
+
+ calculate_urb_setup(devinfo, key, prog_data, shader);
+ brw_compute_flat_inputs(prog_data, shader);
+}
+
/**
* Pre-gen6, the register file of the EUs was shared between threads,
* and each thread used some subset allocated on a 16-register block
void *mem_ctx,
const struct brw_wm_prog_key *key,
struct brw_wm_prog_data *prog_data,
- nir_shader *shader,
+ nir_shader *nir,
int shader_time_index8, int shader_time_index16,
int shader_time_index32, bool allow_spilling,
bool use_rep_send, struct brw_vue_map *vue_map,
const struct gen_device_info *devinfo = compiler->devinfo;
const unsigned max_subgroup_size = compiler->devinfo->gen >= 6 ? 32 : 16;
- brw_nir_apply_key(shader, compiler, &key->base, max_subgroup_size, true);
- brw_nir_lower_fs_inputs(shader, devinfo, key);
- brw_nir_lower_fs_outputs(shader);
+ brw_nir_apply_key(nir, compiler, &key->base, max_subgroup_size, true);
+ brw_nir_lower_fs_inputs(nir, devinfo, key);
+ brw_nir_lower_fs_outputs(nir);
if (devinfo->gen < 6)
- brw_setup_vue_interpolation(vue_map, shader, prog_data);
+ brw_setup_vue_interpolation(vue_map, nir, prog_data);
/* From the SKL PRM, Volume 7, "Alpha Coverage":
* "If Pixel Shader outputs oMask, AlphaToCoverage is disabled in
* offset to determine render target 0 store instruction in
* emit_alpha_to_coverage pass.
*/
- NIR_PASS_V(shader, nir_opt_constant_folding);
- NIR_PASS_V(shader, brw_nir_lower_alpha_to_coverage);
+ NIR_PASS_V(nir, nir_opt_constant_folding);
+ NIR_PASS_V(nir, brw_nir_lower_alpha_to_coverage);
}
if (!key->multisample_fbo)
- NIR_PASS_V(shader, demote_sample_qualifiers);
- NIR_PASS_V(shader, move_interpolation_to_top);
- brw_postprocess_nir(shader, compiler, true);
+ NIR_PASS_V(nir, brw_nir_demote_sample_qualifiers);
+ NIR_PASS_V(nir, brw_nir_move_interpolation_to_top);
+ brw_postprocess_nir(nir, compiler, true);
- /* key->alpha_test_func means simulating alpha testing via discards,
- * so the shader definitely kills pixels.
- */
- prog_data->uses_kill = shader->info.fs.uses_discard ||
- key->alpha_test_func;
- prog_data->uses_omask = key->multisample_fbo &&
- shader->info.outputs_written & BITFIELD64_BIT(FRAG_RESULT_SAMPLE_MASK);
- prog_data->computed_depth_mode = computed_depth_mode(shader);
- prog_data->computed_stencil =
- shader->info.outputs_written & BITFIELD64_BIT(FRAG_RESULT_STENCIL);
-
- prog_data->persample_dispatch =
- key->multisample_fbo &&
- (key->persample_interp ||
- (shader->info.system_values_read & (SYSTEM_BIT_SAMPLE_ID |
- SYSTEM_BIT_SAMPLE_POS)) ||
- shader->info.fs.uses_sample_qualifier ||
- shader->info.outputs_read);
-
- prog_data->has_render_target_reads = shader->info.outputs_read != 0ull;
-
- prog_data->early_fragment_tests = shader->info.fs.early_fragment_tests;
- prog_data->post_depth_coverage = shader->info.fs.post_depth_coverage;
- prog_data->inner_coverage = shader->info.fs.inner_coverage;
-
- prog_data->barycentric_interp_modes =
- brw_compute_barycentric_interp_modes(compiler->devinfo, shader);
-
- calculate_urb_setup(devinfo, key, prog_data, shader);
- brw_compute_flat_inputs(prog_data, shader);
+ brw_nir_populate_wm_prog_data(nir, compiler->devinfo, key, prog_data);
fs_visitor *v8 = NULL, *v16 = NULL, *v32 = NULL;
cfg_t *simd8_cfg = NULL, *simd16_cfg = NULL, *simd32_cfg = NULL;
bool has_spilled = false;
v8 = new fs_visitor(compiler, log_data, mem_ctx, &key->base,
- &prog_data->base, shader, 8, shader_time_index8);
+ &prog_data->base, nir, 8, shader_time_index8);
if (!v8->run_fs(allow_spilling, false /* do_rep_send */)) {
if (error_str)
*error_str = ralloc_strdup(mem_ctx, v8->fail_msg);
likely(!(INTEL_DEBUG & DEBUG_NO16) || use_rep_send)) {
/* Try a SIMD16 compile */
v16 = new fs_visitor(compiler, log_data, mem_ctx, &key->base,
- &prog_data->base, shader, 16, shader_time_index16);
+ &prog_data->base, nir, 16, shader_time_index16);
v16->import_uniforms(v8);
if (!v16->run_fs(allow_spilling, use_rep_send)) {
compiler->shader_perf_log(log_data,
}
}
+ const bool simd16_failed = v16 && !simd16_cfg;
+
/* Currently, the compiler only supports SIMD32 on SNB+ */
if (!has_spilled &&
v8->max_dispatch_width >= 32 && !use_rep_send &&
- devinfo->gen >= 6 && simd16_cfg &&
+ devinfo->gen >= 6 && !simd16_failed &&
!(INTEL_DEBUG & DEBUG_NO32)) {
/* Try a SIMD32 compile */
v32 = new fs_visitor(compiler, log_data, mem_ctx, &key->base,
- &prog_data->base, shader, 32, shader_time_index32);
+ &prog_data->base, nir, 32, shader_time_index32);
v32->import_uniforms(v8);
if (!v32->run_fs(allow_spilling, false)) {
compiler->shader_perf_log(log_data,
if (unlikely(INTEL_DEBUG & DEBUG_WM)) {
g.enable_debug(ralloc_asprintf(mem_ctx, "%s fragment shader %s",
- shader->info.label ?
- shader->info.label : "unnamed",
- shader->info.name));
+ nir->info.label ?
+ nir->info.label : "unnamed",
+ nir->info.name));
}
if (simd8_cfg) {
stats = stats ? stats + 1 : NULL;
}
+ g.add_const_data(nir->constant_data, nir->constant_data_size);
+
delete v8;
delete v16;
delete v32;
fs_reg *
fs_visitor::emit_cs_work_group_id_setup()
{
- assert(stage == MESA_SHADER_COMPUTE);
+ assert(stage == MESA_SHADER_COMPUTE || stage == MESA_SHADER_KERNEL);
fs_reg *reg = new(this->mem_ctx) fs_reg(vgrf(glsl_type::uvec3_type));
void *mem_ctx,
const struct brw_cs_prog_key *key,
struct brw_cs_prog_data *prog_data,
- const nir_shader *src_shader,
+ const nir_shader *nir,
int shader_time_index,
struct brw_compile_stats *stats,
char **error_str)
{
- prog_data->base.total_shared = src_shader->info.cs.shared_size;
- prog_data->slm_size = src_shader->num_shared;
+ prog_data->base.total_shared = nir->info.cs.shared_size;
+ prog_data->slm_size = nir->shared_size;
/* Generate code for all the possible SIMD variants. */
bool generate_all;
unsigned min_dispatch_width;
unsigned max_dispatch_width;
- if (src_shader->info.cs.local_size_variable) {
+ if (nir->info.cs.local_size_variable) {
generate_all = true;
min_dispatch_width = 8;
max_dispatch_width = 32;
} else {
generate_all = false;
- prog_data->local_size[0] = src_shader->info.cs.local_size[0];
- prog_data->local_size[1] = src_shader->info.cs.local_size[1];
- prog_data->local_size[2] = src_shader->info.cs.local_size[2];
+ prog_data->local_size[0] = nir->info.cs.local_size[0];
+ prog_data->local_size[1] = nir->info.cs.local_size[1];
+ prog_data->local_size[2] = nir->info.cs.local_size[2];
unsigned local_workgroup_size = prog_data->local_size[0] *
prog_data->local_size[1] *
prog_data->local_size[2];
if (likely(!(INTEL_DEBUG & DEBUG_NO8)) &&
min_dispatch_width <= 8 && max_dispatch_width >= 8) {
nir_shader *nir8 = compile_cs_to_nir(compiler, mem_ctx, key,
- src_shader, 8);
+ nir, 8);
v8 = new fs_visitor(compiler, log_data, mem_ctx, &key->base,
&prog_data->base,
nir8, 8, shader_time_index);
min_dispatch_width <= 16 && max_dispatch_width >= 16) {
/* Try a SIMD16 compile */
nir_shader *nir16 = compile_cs_to_nir(compiler, mem_ctx, key,
- src_shader, 16);
+ nir, 16);
v16 = new fs_visitor(compiler, log_data, mem_ctx, &key->base,
&prog_data->base,
nir16, 16, shader_time_index);
*
* TODO: Use performance_analysis and drop this boolean.
*/
- const bool needs_32 = min_dispatch_width > 16 ||
+ const bool needs_32 = v == NULL ||
(INTEL_DEBUG & DEBUG_DO32) ||
generate_all;
min_dispatch_width <= 32 && max_dispatch_width >= 32) {
/* Try a SIMD32 compile */
nir_shader *nir32 = compile_cs_to_nir(compiler, mem_ctx, key,
- src_shader, 32);
+ nir, 32);
v32 = new fs_visitor(compiler, log_data, mem_ctx, &key->base,
&prog_data->base,
nir32, 32, shader_time_index);
v->runtime_check_aads_emit, MESA_SHADER_COMPUTE);
if (INTEL_DEBUG & DEBUG_CS) {
char *name = ralloc_asprintf(mem_ctx, "%s compute shader %s",
- src_shader->info.label ?
- src_shader->info.label : "unnamed",
- src_shader->info.name);
+ nir->info.label ?
+ nir->info.label : "unnamed",
+ nir->info.name);
g.enable_debug(name);
}
v->performance_analysis.require(), stats);
}
+ g.add_const_data(nir->constant_data, nir->constant_data_size);
+
ret = g.get_assembly();
delete v8;