* be any component of a vector, and then we load 4 contiguous
* components starting from that.
*
- * We break down the const_offset to a portion added to the variable
- * offset and a portion done using reg_offset, which means that if you
- * have GLSL using something like "uniform vec4 a[20]; gl_FragColor =
- * a[i]", we'll temporarily generate 4 vec4 loads from offset i * 4, and
- * CSE can later notice that those loads are all the same and eliminate
- * the redundant ones.
+ * We break down the const_offset to a portion added to the variable offset
+ * and a portion done using fs_reg::offset, which means that if you have
+ * GLSL using something like "uniform vec4 a[20]; gl_FragColor = a[i]",
+ * we'll temporarily generate 4 vec4 loads from offset i * 4, and CSE can
+ * later notice that those loads are all the same and eliminate the
+ * redundant ones.
*/
fs_reg vec4_offset = vgrf(glsl_type::uint_type);
bld.ADD(vec4_offset, varying_offset, brw_imm_ud(const_offset & ~0xf));
fs_reg vec4_result = bld.vgrf(BRW_REGISTER_TYPE_F, 4);
fs_inst *inst = bld.emit(FS_OPCODE_VARYING_PULL_CONSTANT_LOAD_LOGICAL,
vec4_result, surf_index, vec4_offset);
- inst->size_written = 4 * bld.dispatch_width() / 8 * REG_SIZE;
+ inst->size_written = 4 * vec4_result.component_size(inst->exec_size);
if (type_sz(dst.type) == 8) {
shuffle_32bit_load_result_to_64bit_data(
offset == inst->offset);
}
-bool
-fs_inst::overwrites_reg(const fs_reg ®) const
-{
- return reg.in_range(dst, DIV_ROUND_UP(size_written, REG_SIZE));
-}
-
bool
fs_inst::is_send_from_grf() const
{
return false;
fs_reg reg = this->src[0];
- if (reg.file != VGRF || reg.offset / REG_SIZE != 0 || reg.stride == 0)
+ if (reg.file != VGRF || reg.offset != 0 || reg.stride != 1)
return false;
if (grf_alloc.sizes[reg.nr] * REG_SIZE != this->size_written)
stride == r.stride);
}
-fs_reg &
-fs_reg::set_smear(unsigned subreg)
-{
- assert(file != ARF && file != FIXED_GRF && file != IMM);
- offset = ROUND_DOWN_TO(offset, REG_SIZE) + subreg * type_sz(type);
- stride = 0;
- return *this;
-}
-
bool
fs_reg::is_contiguous() const
{
return 0;
}
-/**
- * Returns the number of scalar components needed to store type, assuming
- * that vectors are padded out to vec4.
- *
- * This has the packing rules of type_size_vec4(), but counts components
- * similar to type_size_scalar().
- */
-extern "C" int
-type_size_vec4_times_4(const struct glsl_type *type)
-{
- return 4 * type_size_vec4(type);
-}
-
-/* Attribute arrays are loaded as one vec4 per element (or matrix column),
- * except for double-precision types, which are loaded as one dvec4.
- */
-extern "C" int
-type_size_vs_input(const struct glsl_type *type)
-{
- if (type->is_double()) {
- return type_size_dvec4(type);
- } else {
- return type_size_vec4(type);
- }
-}
-
/**
* Create a MOV to read the timestamp register.
*
void
fs_visitor::emit_shader_time_begin()
{
- shader_start_time = get_timestamp(bld.annotate("shader time start"));
-
/* We want only the low 32 bits of the timestamp. Since it's running
* at the GPU clock rate of ~1.2ghz, it will roll over every ~3 seconds,
* which is plenty of time for our purposes. It is identical across the
* EUs, but since it's tracking GPU core speed it will increment at a
* varying rate as render P-states change.
*/
- shader_start_time.set_smear(0);
+ shader_start_time = component(
+ get_timestamp(bld.annotate("shader time start")), 0);
}
void
assert(end && ((fs_inst *) end)->eot);
const fs_builder ibld = bld.annotate("shader time end")
.exec_all().at(NULL, end);
-
- fs_reg shader_end_time = get_timestamp(ibld);
+ const fs_reg timestamp = get_timestamp(ibld);
/* We only use the low 32 bits of the timestamp - see
* emit_shader_time_begin()).
* else that might disrupt timing) by setting smear to 2 and checking if
* that field is != 0.
*/
- shader_end_time.set_smear(0);
+ const fs_reg shader_end_time = component(timestamp, 0);
/* Check that there weren't any timestamp reset events (assuming these
* were the only two timestamp reads that happened).
*/
- fs_reg reset = shader_end_time;
- reset.set_smear(2);
+ const fs_reg reset = component(timestamp, 2);
set_condmod(BRW_CONDITIONAL_Z,
ibld.AND(ibld.null_reg_ud(), reset, brw_imm_ud(1u)));
ibld.IF(BRW_PREDICATE_NORMAL);
fs_reg start = shader_start_time;
start.negate = true;
- fs_reg diff = fs_reg(VGRF, alloc.allocate(1), BRW_REGISTER_TYPE_UD);
- diff.set_smear(0);
-
+ const fs_reg diff = component(fs_reg(VGRF, alloc.allocate(1),
+ BRW_REGISTER_TYPE_UD),
+ 0);
const fs_builder cbld = ibld.group(1, 0);
cbld.group(1, 0).ADD(diff, start, shader_end_time);
opcode == SHADER_OPCODE_TXD_LOGICAL)
return src[TEX_LOGICAL_SRC_GRAD_COMPONENTS].ud;
/* Texture offset. */
- else if (i == TEX_LOGICAL_SRC_OFFSET_VALUE)
+ else if (i == TEX_LOGICAL_SRC_TG4_OFFSET)
return 2;
/* MCS */
else if (i == TEX_LOGICAL_SRC_MCS && opcode == SHADER_OPCODE_TXF_CMS_W_LOGICAL)
}
}
-int
+unsigned
fs_inst::size_read(int arg) const
{
switch (opcode) {
case FS_OPCODE_LINTERP:
if (arg == 1)
- return REG_SIZE;
+ return 16;
break;
case SHADER_OPCODE_LOAD_PAYLOAD:
case SHADER_OPCODE_MOV_INDIRECT:
if (arg == 0) {
assert(src[2].file == IMM);
- unsigned region_length = src[2].ud;
-
- if (src[0].file == UNIFORM) {
- assert(region_length % 4 == 0);
- return region_length;
- } else if (src[0].file == FIXED_GRF) {
- /* If the start of the region is not register aligned, then
- * there's some portion of the register that's technically
- * unread at the beginning.
- *
- * However, the register allocator works in terms of whole
- * registers, and does not use subnr. It assumes that the
- * read starts at the beginning of the register, and extends
- * regs_read() whole registers beyond that.
- *
- * To compensate, we extend the region length to include this
- * unread portion at the beginning.
- */
- if (src[0].subnr)
- region_length += src[0].subnr;
-
- return region_length;
- } else {
- assert(!"Invalid register file");
- }
+ return src[2].ud;
}
break;
fs_visitor::compute_sample_position(fs_reg dst, fs_reg int_sample_pos)
{
assert(stage == MESA_SHADER_FRAGMENT);
- brw_wm_prog_data *wm_prog_data = (brw_wm_prog_data *) this->prog_data;
+ struct brw_wm_prog_data *wm_prog_data = brw_wm_prog_data(this->prog_data);
assert(dst.type == BRW_REGISTER_TYPE_F);
if (wm_prog_data->persample_dispatch) {
brw_imm_v(0x44440000));
abld.AND(*reg, tmp, brw_imm_w(0xf));
} else {
- fs_reg t1(VGRF, alloc.allocate(1), BRW_REGISTER_TYPE_D);
- t1.set_smear(0);
- fs_reg t2(VGRF, alloc.allocate(1), BRW_REGISTER_TYPE_W);
+ const fs_reg t1 = component(fs_reg(VGRF, alloc.allocate(1),
+ BRW_REGISTER_TYPE_D), 0);
+ const fs_reg t2(VGRF, alloc.allocate(1), BRW_REGISTER_TYPE_W);
/* The PS will be run in MSDISPMODE_PERSAMPLE. For example with
* 8x multisampling, subspan 0 will represent sample N (where N
fs_visitor::emit_samplemaskin_setup()
{
assert(stage == MESA_SHADER_FRAGMENT);
- brw_wm_prog_data *wm_prog_data = (brw_wm_prog_data *) this->prog_data;
+ struct brw_wm_prog_data *wm_prog_data = brw_wm_prog_data(this->prog_data);
assert(devinfo->gen >= 6);
fs_reg *reg = new(this->mem_ctx) fs_reg(vgrf(glsl_type::int_type));
void
fs_visitor::emit_discard_jump()
{
- assert(((brw_wm_prog_data*) this->prog_data)->uses_kill);
+ assert(brw_wm_prog_data(this->prog_data)->uses_kill);
/* For performance, after a discard, jump to the end of the
* shader if all relevant channels have been discarded.
{
assert(stage == MESA_SHADER_GEOMETRY);
- struct brw_gs_prog_data *gs_prog_data =
- (struct brw_gs_prog_data *) prog_data;
+ struct brw_gs_prog_data *gs_prog_data = brw_gs_prog_data(prog_data);
if (gs_compile->control_data_header_size_bits > 0) {
emit_gs_control_data_bits(this->final_gs_vertex_count);
fs_visitor::calculate_urb_setup()
{
assert(stage == MESA_SHADER_FRAGMENT);
- brw_wm_prog_data *prog_data = (brw_wm_prog_data*) this->prog_data;
+ struct brw_wm_prog_data *prog_data = brw_wm_prog_data(this->prog_data);
brw_wm_prog_key *key = (brw_wm_prog_key*) this->key;
memset(prog_data->urb_setup, -1,
int urb_next = 0;
/* Figure out where each of the incoming setup attributes lands. */
if (devinfo->gen >= 6) {
- if (_mesa_bitcount_64(nir->info.inputs_read &
+ if (_mesa_bitcount_64(nir->info->inputs_read &
BRW_FS_VARYING_INPUT_MASK) <= 16) {
/* The SF/SBE pipeline stage can do arbitrary rearrangement of the
* first 16 varying inputs, so we can put them wherever we want.
* a different vertex (or geometry) shader.
*/
for (unsigned int i = 0; i < VARYING_SLOT_MAX; i++) {
- if (nir->info.inputs_read & BRW_FS_VARYING_INPUT_MASK &
+ if (nir->info->inputs_read & BRW_FS_VARYING_INPUT_MASK &
BITFIELD64_BIT(i)) {
prog_data->urb_setup[i] = urb_next++;
}
}
} else {
bool include_vue_header =
- nir->info.inputs_read & (VARYING_BIT_LAYER | VARYING_BIT_VIEWPORT);
+ nir->info->inputs_read & (VARYING_BIT_LAYER | VARYING_BIT_VIEWPORT);
/* We have enough input varyings that the SF/SBE pipeline stage can't
* arbitrarily rearrange them to suit our whim; we have to put them
struct brw_vue_map prev_stage_vue_map;
brw_compute_vue_map(devinfo, &prev_stage_vue_map,
key->input_slots_valid,
- nir->info.separate_shader);
+ nir->info->separate_shader);
int first_slot =
include_vue_header ? 0 : 2 * BRW_SF_URB_ENTRY_READ_OFFSET;
slot++) {
int varying = prev_stage_vue_map.slot_to_varying[slot];
if (varying != BRW_VARYING_SLOT_PAD &&
- (nir->info.inputs_read & BRW_FS_VARYING_INPUT_MASK &
+ (nir->info->inputs_read & BRW_FS_VARYING_INPUT_MASK &
BITFIELD64_BIT(varying))) {
prog_data->urb_setup[varying] = slot - first_slot;
}
*
* See compile_sf_prog() for more info.
*/
- if (nir->info.inputs_read & BITFIELD64_BIT(VARYING_SLOT_PNTC))
+ if (nir->info->inputs_read & BITFIELD64_BIT(VARYING_SLOT_PNTC))
prog_data->urb_setup[VARYING_SLOT_PNTC] = urb_next++;
}
fs_visitor::assign_urb_setup()
{
assert(stage == MESA_SHADER_FRAGMENT);
- brw_wm_prog_data *prog_data = (brw_wm_prog_data*) this->prog_data;
+ struct brw_wm_prog_data *prog_data = brw_wm_prog_data(this->prog_data);
int urb_start = payload.num_regs + prog_data->base.curb_read_length;
void
fs_visitor::assign_vs_urb_setup()
{
- brw_vs_prog_data *vs_prog_data = (brw_vs_prog_data *) prog_data;
+ struct brw_vs_prog_data *vs_prog_data = brw_vs_prog_data(prog_data);
assert(stage == MESA_SHADER_VERTEX);
{
assert(stage == MESA_SHADER_TESS_EVAL);
- brw_vue_prog_data *vue_prog_data = (brw_vue_prog_data *) prog_data;
+ struct brw_vue_prog_data *vue_prog_data = brw_vue_prog_data(prog_data);
first_non_payload_grf += 8 * vue_prog_data->urb_read_length;
{
assert(stage == MESA_SHADER_GEOMETRY);
- brw_vue_prog_data *vue_prog_data = (brw_vue_prog_data *) prog_data;
+ struct brw_vue_prog_data *vue_prog_data = brw_vue_prog_data(prog_data);
first_non_payload_grf +=
- 8 * vue_prog_data->urb_read_length * nir->info.gs.vertices_in;
+ 8 * vue_prog_data->urb_read_length * nir->info->gs.vertices_in;
foreach_block_and_inst(block, fs_inst, inst, cfg) {
/* Rewrite all ATTR file references to GRFs. */
void
fs_visitor::split_virtual_grfs()
{
+ /* Compact the register file so we eliminate dead vgrfs. This
+ * only defines split points for live registers, so if we have
+ * too large dead registers they will hit assertions later.
+ */
+ compact_virtual_grfs();
+
int num_vars = this->alloc.count;
/* Count the total number of registers */
int thread_local_id_index =
(stage == MESA_SHADER_COMPUTE) ?
- ((brw_cs_prog_data*)stage_prog_data)->thread_local_id_index : -1;
+ brw_cs_prog_data(stage_prog_data)->thread_local_id_index : -1;
/* First, we walk through the instructions and do two things:
*
stage_prog_data->nr_params = num_push_constants;
stage_prog_data->nr_pull_params = num_pull_constants;
- /* Up until now, the param[] array has been indexed by reg + reg_offset
+ /* Up until now, the param[] array has been indexed by reg + offset
* of UNIFORM registers. Move pull constants into pull_param[] and
* condense param[] to only contain the uniforms we chose to push.
*
ralloc_free(param);
if (stage == MESA_SHADER_COMPUTE)
- ((brw_cs_prog_data*)stage_prog_data)->thread_local_id_index =
+ brw_cs_prog_data(stage_prog_data)->thread_local_id_index =
new_thread_local_id_index;
}
if (pull_index == -1)
continue;
- const unsigned index = stage_prog_data->binding_table.pull_constants_start;
- fs_reg dst;
-
- if (type_sz(inst->src[i].type) <= 4)
- dst = vgrf(glsl_type::float_type);
- else
- dst = vgrf(glsl_type::double_type);
-
assert(inst->src[i].stride == 0);
- const fs_builder ubld = ibld.exec_all().group(8, 0);
- struct brw_reg offset = brw_imm_ud((unsigned)(pull_index * 4) & ~15);
+ const unsigned index = stage_prog_data->binding_table.pull_constants_start;
+ const unsigned block_sz = 64; /* Fetch one cacheline at a time. */
+ const fs_builder ubld = ibld.exec_all().group(block_sz / 4, 0);
+ const fs_reg dst = ubld.vgrf(BRW_REGISTER_TYPE_UD);
+ const unsigned base = pull_index * 4;
+
ubld.emit(FS_OPCODE_UNIFORM_PULL_CONSTANT_LOAD,
- dst, brw_imm_ud(index), offset);
+ dst, brw_imm_ud(index), brw_imm_ud(base & ~(block_sz - 1)));
/* Rewrite the instruction to use the temporary VGRF. */
inst->src[i].file = VGRF;
inst->src[i].nr = dst.nr;
- inst->src[i].offset %= 4;
- inst->src[i].set_smear((pull_index & 3) * 4 /
- type_sz(inst->src[i].type));
+ inst->src[i].offset = (base & (block_sz - 1)) +
+ inst->src[i].offset % 4;
brw_mark_surface_used(prog_data, index);
}
/**
* Compute a bitmask with GRF granularity with a bit set for each GRF starting
- * from \p r which overlaps the region starting at \p r and spanning \p n GRF
- * units.
+ * from \p r.offset which overlaps the region starting at \p s.offset and
+ * spanning \p ds bytes.
*/
static inline unsigned
-mask_relative_to(const fs_reg &r, const fs_reg &s, unsigned n)
+mask_relative_to(const fs_reg &r, const fs_reg &s, unsigned ds)
{
- const int rel_offset = (reg_offset(s) - reg_offset(r)) / REG_SIZE;
+ const int rel_offset = reg_offset(s) - reg_offset(r);
+ const int shift = rel_offset / REG_SIZE;
+ const unsigned n = DIV_ROUND_UP(rel_offset % REG_SIZE + ds, REG_SIZE);
assert(reg_space(r) == reg_space(s) &&
- rel_offset >= 0 && rel_offset < int(8 * sizeof(unsigned)));
- return ((1 << n) - 1) << rel_offset;
+ shift >= 0 && shift < int(8 * sizeof(unsigned)));
+ return ((1 << n) - 1) << shift;
}
bool
* would need us to understand coalescing out more than one MOV at
* a time.
*/
- if (scan_inst->dst.offset / REG_SIZE < inst->src[0].offset / REG_SIZE ||
- scan_inst->dst.offset / REG_SIZE + DIV_ROUND_UP(scan_inst->size_written, REG_SIZE) >
- inst->src[0].offset / REG_SIZE + DIV_ROUND_UP(inst->size_read(0), REG_SIZE))
+ if (!region_contained_in(scan_inst->dst, scan_inst->size_written,
+ inst->src[0], inst->size_read(0)))
break;
/* SEND instructions can't have MRF as a destination. */
/* Clear the bits for any registers this instruction overwrites. */
regs_left &= ~mask_relative_to(
- inst->src[0], scan_inst->dst, DIV_ROUND_UP(scan_inst->size_written,
- REG_SIZE));
+ inst->src[0], scan_inst->dst, scan_inst->size_written);
if (!regs_left)
break;
}
inst->src[0], inst->size_read(0))) {
/* Clear the bits for any registers this instruction overwrites. */
regs_left &= ~mask_relative_to(
- inst->src[0], scan_inst->dst, DIV_ROUND_UP(scan_inst->size_written,
- REG_SIZE));
+ inst->src[0], scan_inst->dst, scan_inst->size_written);
- const unsigned rel_offset = (reg_offset(scan_inst->dst) -
- reg_offset(inst->src[0])) / REG_SIZE;
+ const unsigned rel_offset = reg_offset(scan_inst->dst) -
+ reg_offset(inst->src[0]);
if (inst->dst.nr & BRW_MRF_COMPR4) {
/* Apply the same address transformation done by the hardware
* for COMPR4 MRF writes.
*/
- assert(rel_offset < 2);
- scan_inst->dst.nr = inst->dst.nr + rel_offset * 4;
+ assert(rel_offset < 2 * REG_SIZE);
+ scan_inst->dst.nr = inst->dst.nr + rel_offset / REG_SIZE * 4;
/* Clear the COMPR4 bit if the generating instruction is not
* compressed.
/* Calculate the MRF number the result of this instruction is
* ultimately written to.
*/
- scan_inst->dst.nr = inst->dst.nr + rel_offset;
+ scan_inst->dst.nr = inst->dst.nr + rel_offset / REG_SIZE;
}
scan_inst->dst.file = MRF;
- scan_inst->dst.offset %= REG_SIZE;
+ scan_inst->dst.offset = inst->dst.offset + rel_offset % REG_SIZE;
scan_inst->saturate |= inst->saturate;
if (!regs_left)
break;
bool progress = false;
unsigned depth = 0;
+ if (!brw_stage_has_packed_dispatch(devinfo, stage, stage_prog_data)) {
+ /* The optimization below assumes that channel zero is live on thread
+ * dispatch, which may not be the case if the fixed function dispatches
+ * threads sparsely.
+ */
+ return false;
+ }
+
foreach_block_and_inst_safe(block, fs_inst, inst, cfg) {
switch (inst->opcode) {
case BRW_OPCODE_IF:
bool progress = false;
- /* Note that we're done with register allocation, so GRF fs_regs always
- * have a .reg_offset of 0.
- */
-
foreach_block_and_inst(block, fs_inst, inst, cfg) {
if (inst->mlen != 0 && inst->dst.file == VGRF) {
insert_gen4_pre_send_dependency_workarounds(block, inst);
continue;
if (devinfo->gen >= 7) {
- /* The offset arg is a vec4-aligned immediate byte offset. */
- fs_reg const_offset_reg = inst->src[1];
- assert(const_offset_reg.file == IMM &&
- const_offset_reg.type == BRW_REGISTER_TYPE_UD);
- assert(const_offset_reg.ud % 16 == 0);
-
- fs_reg payload, offset;
- if (devinfo->gen >= 9) {
- /* We have to use a message header on Skylake to get SIMD4x2
- * mode. Reserve space for the register.
- */
- offset = payload = fs_reg(VGRF, alloc.allocate(2));
- offset.offset += REG_SIZE;
- inst->mlen = 2;
- } else {
- offset = payload = fs_reg(VGRF, alloc.allocate(1));
- inst->mlen = 1;
- }
-
- /* This is actually going to be a MOV, but since only the first dword
- * is accessed, we have a special opcode to do just that one. Note
- * that this needs to be an operation that will be considered a def
- * by live variable analysis, or register allocation will explode.
- */
- fs_inst *setup = new(mem_ctx) fs_inst(FS_OPCODE_SET_SIMD4X2_OFFSET,
- 8, offset, const_offset_reg);
- setup->force_writemask_all = true;
+ const fs_builder ubld = fs_builder(this, block, inst).exec_all();
+ const fs_reg payload = ubld.group(8, 0).vgrf(BRW_REGISTER_TYPE_UD);
- setup->ir = inst->ir;
- setup->annotation = inst->annotation;
- inst->insert_before(block, setup);
+ ubld.group(8, 0).MOV(payload,
+ retype(brw_vec8_grf(0, 0), BRW_REGISTER_TYPE_UD));
+ ubld.group(1, 0).MOV(component(payload, 2),
+ brw_imm_ud(inst->src[1].ud / 16));
- /* Similarly, this will only populate the first 4 channels of the
- * result register (since we only use smear values from 0-3), but we
- * don't tell the optimizer.
- */
inst->opcode = FS_OPCODE_UNIFORM_PULL_CONSTANT_LOAD_GEN7;
inst->src[1] = payload;
+ inst->header_size = 1;
+ inst->mlen = 1;
invalidate_live_intervals();
} else {
inst->dst.type);
if (devinfo->gen >= 7) {
- fs_reg src1_0_w = inst->src[1];
- fs_reg src1_1_w = inst->src[1];
-
if (inst->src[1].file == IMM) {
- src1_0_w.ud &= 0xffff;
- src1_1_w.ud >>= 16;
+ ibld.MUL(low, inst->src[0],
+ brw_imm_uw(inst->src[1].ud & 0xffff));
+ ibld.MUL(high, inst->src[0],
+ brw_imm_uw(inst->src[1].ud >> 16));
} else {
- src1_0_w.type = BRW_REGISTER_TYPE_UW;
- if (src1_0_w.stride != 0) {
- assert(src1_0_w.stride == 1);
- src1_0_w.stride = 2;
- }
-
- src1_1_w.type = BRW_REGISTER_TYPE_UW;
- if (src1_1_w.stride != 0) {
- assert(src1_1_w.stride == 1);
- src1_1_w.stride = 2;
- }
- src1_1_w.offset += type_sz(BRW_REGISTER_TYPE_UW);
+ ibld.MUL(low, inst->src[0],
+ subscript(inst->src[1], BRW_REGISTER_TYPE_UW, 0));
+ ibld.MUL(high, inst->src[0],
+ subscript(inst->src[1], BRW_REGISTER_TYPE_UW, 1));
}
- ibld.MUL(low, inst->src[0], src1_0_w);
- ibld.MUL(high, inst->src[0], src1_1_w);
} else {
- fs_reg src0_0_w = inst->src[0];
- fs_reg src0_1_w = inst->src[0];
-
- src0_0_w.type = BRW_REGISTER_TYPE_UW;
- if (src0_0_w.stride != 0) {
- assert(src0_0_w.stride == 1);
- src0_0_w.stride = 2;
- }
-
- src0_1_w.type = BRW_REGISTER_TYPE_UW;
- if (src0_1_w.stride != 0) {
- assert(src0_1_w.stride == 1);
- src0_1_w.stride = 2;
- }
- src0_1_w.offset += type_sz(BRW_REGISTER_TYPE_UW);
-
- ibld.MUL(low, src0_0_w, inst->src[1]);
- ibld.MUL(high, src0_1_w, inst->src[1]);
+ ibld.MUL(low, subscript(inst->src[0], BRW_REGISTER_TYPE_UW, 0),
+ inst->src[1]);
+ ibld.MUL(high, subscript(inst->src[0], BRW_REGISTER_TYPE_UW, 1),
+ inst->src[1]);
}
- fs_reg dst = inst->dst;
- dst.type = BRW_REGISTER_TYPE_UW;
- dst.offset = ROUND_DOWN_TO(dst.offset, REG_SIZE) + 2;
- dst.stride = 2;
-
- high.type = BRW_REGISTER_TYPE_UW;
- high.stride = 2;
-
- low.type = BRW_REGISTER_TYPE_UW;
- low.offset = ROUND_DOWN_TO(low.offset, REG_SIZE) + 2;
- low.stride = 2;
-
- ibld.ADD(dst, low, high);
+ ibld.ADD(subscript(inst->dst, BRW_REGISTER_TYPE_UW, 1),
+ subscript(low, BRW_REGISTER_TYPE_UW, 1),
+ subscript(high, BRW_REGISTER_TYPE_UW, 0));
if (inst->conditional_mod || orig_dst.file == MRF) {
set_condmod(inst->conditional_mod,
static void
lower_fb_write_logical_send(const fs_builder &bld, fs_inst *inst,
- const brw_wm_prog_data *prog_data,
+ const struct brw_wm_prog_data *prog_data,
const brw_wm_prog_key *key,
const fs_visitor::thread_payload &payload)
{
*/
setup_color_payload(bld, key, &sources[length], src0_alpha, 1);
length++;
+ } else if (key->replicate_alpha && inst->target != 0) {
+ /* Handle the case when fragment shader doesn't write to draw buffer
+ * zero. No need to call setup_color_payload() for src0_alpha because
+ * alpha value will be undefined.
+ */
+ length++;
}
setup_color_payload(bld, key, &sources[length], color0, components);
}
if (has_lod) {
- /* Bias/LOD with shadow comparitor is unsupported in SIMD16 -- *Without*
- * shadow comparitor (including RESINFO) it's unsupported in SIMD8 mode.
+ /* Bias/LOD with shadow comparator is unsupported in SIMD16 -- *Without*
+ * shadow comparator (including RESINFO) it's unsupported in SIMD8 mode.
*/
assert(shadow_c.file != BAD_FILE ? bld.dispatch_width() == 8 :
bld.dispatch_width() == 16);
const fs_reg &sample_index,
const fs_reg &surface,
const fs_reg &sampler,
- const fs_reg &offset_value,
unsigned coord_components,
unsigned grad_components)
{
fs_reg msg_coords = message;
unsigned header_size = 0;
- if (offset_value.file != BAD_FILE) {
+ if (inst->offset != 0) {
/* The offsets set up by the visitor are in the m1 header, so we can't
* go headerless.
*/
const fs_reg &mcs,
const fs_reg &surface,
const fs_reg &sampler,
- const fs_reg &offset_value,
+ const fs_reg &tg4_offset,
unsigned coord_components,
unsigned grad_components)
{
sources[i] = bld.vgrf(BRW_REGISTER_TYPE_F);
if (op == SHADER_OPCODE_TG4 || op == SHADER_OPCODE_TG4_OFFSET ||
- offset_value.file != BAD_FILE || inst->eot ||
+ inst->offset != 0 || inst->eot ||
op == SHADER_OPCODE_SAMPLEINFO ||
is_high_sampler(devinfo, sampler)) {
/* For general texture offsets (no txf workaround), we need a header to
for (unsigned i = 0; i < 2; i++) /* offu, offv */
bld.MOV(retype(sources[length++], BRW_REGISTER_TYPE_D),
- offset(offset_value, bld, i));
+ offset(tg4_offset, bld, i));
if (coord_components == 3) /* r if present */
bld.MOV(sources[length++], offset(coordinate, bld, 2));
const fs_reg &mcs = inst->src[TEX_LOGICAL_SRC_MCS];
const fs_reg &surface = inst->src[TEX_LOGICAL_SRC_SURFACE];
const fs_reg &sampler = inst->src[TEX_LOGICAL_SRC_SAMPLER];
- const fs_reg &offset_value = inst->src[TEX_LOGICAL_SRC_OFFSET_VALUE];
+ const fs_reg &tg4_offset = inst->src[TEX_LOGICAL_SRC_TG4_OFFSET];
assert(inst->src[TEX_LOGICAL_SRC_COORD_COMPONENTS].file == IMM);
const unsigned coord_components = inst->src[TEX_LOGICAL_SRC_COORD_COMPONENTS].ud;
assert(inst->src[TEX_LOGICAL_SRC_GRAD_COMPONENTS].file == IMM);
if (devinfo->gen >= 7) {
lower_sampler_logical_send_gen7(bld, inst, op, coordinate,
shadow_c, lod, lod2, sample_index,
- mcs, surface, sampler, offset_value,
+ mcs, surface, sampler, tg4_offset,
coord_components, grad_components);
} else if (devinfo->gen >= 5) {
lower_sampler_logical_send_gen5(bld, inst, op, coordinate,
shadow_c, lod, lod2, sample_index,
- surface, sampler, offset_value,
+ surface, sampler,
coord_components, grad_components);
} else {
lower_sampler_logical_send_gen4(bld, inst, op, coordinate,
case FS_OPCODE_FB_WRITE_LOGICAL:
assert(stage == MESA_SHADER_FRAGMENT);
lower_fb_write_logical_send(ibld, inst,
- (const brw_wm_prog_data *)prog_data,
+ brw_wm_prog_data(prog_data),
(const brw_wm_prog_key *)key,
payload);
break;
*/
if (devinfo->gen < 8) {
for (unsigned i = 0; i < inst->sources; i++) {
- if (DIV_ROUND_UP(inst->size_written, REG_SIZE) == 2 &&
- inst->size_read(i) != 0 && DIV_ROUND_UP(inst->size_read(i), REG_SIZE) != 2 &&
+ if (inst->size_written > REG_SIZE &&
+ inst->size_read(i) != 0 && inst->size_read(i) <= REG_SIZE &&
!is_uniform(inst->src[i]) &&
!(type_sz(inst->dst.type) == 4 && inst->dst.stride == 1 &&
type_sz(inst->src[i].type) == 2 && inst->src[i].stride == 1)) {
inst->components_read(TEX_LOGICAL_SRC_LOD2) +
inst->components_read(TEX_LOGICAL_SRC_SAMPLE_INDEX) +
(inst->opcode == SHADER_OPCODE_TG4_OFFSET_LOGICAL ?
- inst->components_read(TEX_LOGICAL_SRC_OFFSET_VALUE) : 0) +
+ inst->components_read(TEX_LOGICAL_SRC_TG4_OFFSET) : 0) +
inst->components_read(TEX_LOGICAL_SRC_MCS);
/* SIMD16 messages with more than five arguments exceed the maximum message
switch (inst->dst.file) {
case VGRF:
fprintf(file, "vgrf%d", inst->dst.nr);
- if (alloc.sizes[inst->dst.nr] * REG_SIZE != inst->size_written ||
- inst->dst.offset % REG_SIZE)
- fprintf(file, "+%d.%d",
- inst->dst.offset / REG_SIZE, inst->dst.offset % REG_SIZE);
break;
case FIXED_GRF:
fprintf(file, "g%d", inst->dst.nr);
fprintf(file, "(null)");
break;
case UNIFORM:
- fprintf(file, "***u%d***", inst->dst.nr + inst->dst.offset / 4);
+ fprintf(file, "***u%d***", inst->dst.nr);
break;
case ATTR:
- fprintf(file, "***attr%d***", inst->dst.nr + inst->dst.offset / REG_SIZE);
+ fprintf(file, "***attr%d***", inst->dst.nr);
break;
case ARF:
switch (inst->dst.nr) {
fprintf(file, "arf%d.%d", inst->dst.nr & 0xf, inst->dst.subnr);
break;
}
- if (inst->dst.subnr)
- fprintf(file, "+%d", inst->dst.subnr);
break;
case IMM:
unreachable("not reached");
}
+
+ if (inst->dst.offset ||
+ (inst->dst.file == VGRF &&
+ alloc.sizes[inst->dst.nr] * REG_SIZE != inst->size_written)) {
+ const unsigned reg_size = (inst->dst.file == UNIFORM ? 4 : REG_SIZE);
+ fprintf(file, "+%d.%d", inst->dst.offset / reg_size,
+ inst->dst.offset % reg_size);
+ }
+
if (inst->dst.stride != 1)
fprintf(file, "<%u>", inst->dst.stride);
fprintf(file, ":%s, ", brw_reg_type_letters(inst->dst.type));
switch (inst->src[i].file) {
case VGRF:
fprintf(file, "vgrf%d", inst->src[i].nr);
- if (alloc.sizes[inst->src[i].nr] * REG_SIZE != inst->size_read(i) ||
- inst->src[i].offset % REG_SIZE != 0)
- fprintf(file, "+%d.%d", inst->src[i].offset / REG_SIZE,
- inst->src[i].offset % REG_SIZE);
break;
case FIXED_GRF:
fprintf(file, "g%d", inst->src[i].nr);
fprintf(file, "***m%d***", inst->src[i].nr);
break;
case ATTR:
- fprintf(file, "attr%d+%d", inst->src[i].nr, inst->src[i].offset / REG_SIZE);
+ fprintf(file, "attr%d", inst->src[i].nr);
break;
case UNIFORM:
- fprintf(file, "u%d", inst->src[i].nr + inst->src[i].offset / 4);
- if (inst->src[i].offset % 4 != 0) {
- fprintf(file, "+%d.%d", inst->src[i].offset / 4,
- inst->src[i].offset % 4);
- }
+ fprintf(file, "u%d", inst->src[i].nr);
break;
case BAD_FILE:
fprintf(file, "(null)");
fprintf(file, "arf%d.%d", inst->src[i].nr & 0xf, inst->src[i].subnr);
break;
}
- if (inst->src[i].subnr)
- fprintf(file, "+%d", inst->src[i].subnr);
break;
}
+
+ if (inst->src[i].offset ||
+ (inst->src[i].file == VGRF &&
+ alloc.sizes[inst->src[i].nr] * REG_SIZE != inst->size_read(i))) {
+ const unsigned reg_size = (inst->src[i].file == UNIFORM ? 4 : REG_SIZE);
+ fprintf(file, "+%d.%d", inst->src[i].offset / reg_size,
+ inst->src[i].offset % reg_size);
+ }
+
if (inst->src[i].abs)
fprintf(file, "|");
fs_visitor::setup_fs_payload_gen6()
{
assert(stage == MESA_SHADER_FRAGMENT);
- brw_wm_prog_data *prog_data = (brw_wm_prog_data*) this->prog_data;
-
- unsigned barycentric_interp_modes =
- (stage == MESA_SHADER_FRAGMENT) ?
- ((brw_wm_prog_data*) this->prog_data)->barycentric_interp_modes : 0;
+ struct brw_wm_prog_data *prog_data = brw_wm_prog_data(this->prog_data);
assert(devinfo->gen >= 6);
* Mode" bits in WM_STATE.
*/
for (int i = 0; i < BRW_BARYCENTRIC_MODE_COUNT; ++i) {
- if (barycentric_interp_modes & (1 << i)) {
+ if (prog_data->barycentric_interp_modes & (1 << i)) {
payload.barycentric_coord_reg[i] = payload.num_regs;
payload.num_regs += 2;
if (dispatch_width == 16) {
/* R27: interpolated depth if uses source depth */
prog_data->uses_src_depth =
- (nir->info.inputs_read & (1 << VARYING_SLOT_POS)) != 0;
+ (nir->info->inputs_read & (1 << VARYING_SLOT_POS)) != 0;
if (prog_data->uses_src_depth) {
payload.source_depth_reg = payload.num_regs;
payload.num_regs++;
/* R29: interpolated W set if GEN6_WM_USES_SOURCE_W. */
prog_data->uses_src_w =
- (nir->info.inputs_read & (1 << VARYING_SLOT_POS)) != 0;
+ (nir->info->inputs_read & (1 << VARYING_SLOT_POS)) != 0;
if (prog_data->uses_src_w) {
payload.source_w_reg = payload.num_regs;
payload.num_regs++;
/* R31: MSAA position offsets. */
if (prog_data->persample_dispatch &&
- (nir->info.system_values_read & SYSTEM_BIT_SAMPLE_POS)) {
+ (nir->info->system_values_read & SYSTEM_BIT_SAMPLE_POS)) {
/* From the Ivy Bridge PRM documentation for 3DSTATE_PS:
*
* "MSDISPMODE_PERSAMPLE is required in order to select
/* R32: MSAA input coverage mask */
prog_data->uses_sample_mask =
- (nir->info.system_values_read & SYSTEM_BIT_SAMPLE_MASK_IN) != 0;
+ (nir->info->system_values_read & SYSTEM_BIT_SAMPLE_MASK_IN) != 0;
if (prog_data->uses_sample_mask) {
assert(devinfo->gen >= 7);
payload.sample_mask_in_reg = payload.num_regs;
/* R34-: bary for 32-pixel. */
/* R58-59: interp W for 32-pixel. */
- if (nir->info.outputs_written & BITFIELD64_BIT(FRAG_RESULT_DEPTH)) {
+ if (nir->info->outputs_written & BITFIELD64_BIT(FRAG_RESULT_DEPTH)) {
source_depth_to_render_target = true;
}
}
{
assert(stage == MESA_SHADER_GEOMETRY);
- struct brw_gs_prog_data *gs_prog_data =
- (struct brw_gs_prog_data *) prog_data;
- struct brw_vue_prog_data *vue_prog_data =
- (struct brw_vue_prog_data *) prog_data;
+ struct brw_gs_prog_data *gs_prog_data = brw_gs_prog_data(prog_data);
+ struct brw_vue_prog_data *vue_prog_data = brw_vue_prog_data(prog_data);
/* R0: thread header, R1: output URB handles */
payload.num_regs = 2;
* Note that the GS reads <URB Read Length> HWords for every vertex - so we
* have to multiply by VerticesIn to obtain the total storage requirement.
*/
- if (8 * vue_prog_data->urb_read_length * nir->info.gs.vertices_in >
+ if (8 * vue_prog_data->urb_read_length * nir->info->gs.vertices_in >
max_push_components || gs_prog_data->invocations > 1) {
gs_prog_data->base.include_vue_handles = true;
/* R3..RN: ICP Handles for each incoming vertex (when using pull model) */
- payload.num_regs += nir->info.gs.vertices_in;
+ payload.num_regs += nir->info->gs.vertices_in;
vue_prog_data->urb_read_length =
- ROUND_DOWN_TO(max_push_components / nir->info.gs.vertices_in, 8) / 8;
+ ROUND_DOWN_TO(max_push_components / nir->info->gs.vertices_in, 8) / 8;
}
}
if (unlikely(INTEL_DEBUG & DEBUG_OPTIMIZER) && this_progress) { \
char filename[64]; \
snprintf(filename, 64, "%s%d-%s-%02d-%02d-" #pass, \
- stage_abbrev, dispatch_width, nir->info.name, iteration, pass_num); \
+ stage_abbrev, dispatch_width, nir->info->name, iteration, pass_num); \
\
backend_shader::dump_instructions(filename); \
} \
if (unlikely(INTEL_DEBUG & DEBUG_OPTIMIZER)) {
char filename[64];
snprintf(filename, 64, "%s%d-%s-00-00-start",
- stage_abbrev, dispatch_width, nir->info.name);
+ stage_abbrev, dispatch_width, nir->info->name);
backend_shader::dump_instructions(filename);
}
OPT(opt_algebraic);
OPT(opt_cse);
- OPT(opt_copy_propagate);
+ OPT(opt_copy_propagation);
OPT(opt_predicated_break, this);
OPT(opt_cmod_propagation);
OPT(dead_code_eliminate);
}
if (OPT(lower_d2x)) {
- OPT(opt_copy_propagate);
+ OPT(opt_copy_propagation);
OPT(dead_code_eliminate);
}
OPT(lower_logical_sends);
if (progress) {
- OPT(opt_copy_propagate);
+ OPT(opt_copy_propagation);
/* Only run after logical send lowering because it's easier to implement
* in terms of physical sends.
*/
if (OPT(opt_zero_samples))
- OPT(opt_copy_propagate);
+ OPT(opt_copy_propagation);
/* Run after logical send lowering to give it a chance to CSE the
* LOAD_PAYLOAD instructions created to construct the payloads of
* e.g. texturing messages in cases where it wasn't possible to CSE the
if (devinfo->gen <= 5 && OPT(lower_minmax)) {
OPT(opt_cmod_propagation);
OPT(opt_cse);
- OPT(opt_copy_propagate);
+ OPT(opt_copy_propagation);
OPT(dead_code_eliminate);
}
schedule_instructions(SCHEDULE_POST);
if (last_scratch > 0) {
- unsigned max_scratch_size = 2 * 1024 * 1024;
+ MAYBE_UNUSED unsigned max_scratch_size = 2 * 1024 * 1024;
prog_data->total_scratch = brw_get_scratch_size(last_scratch);
{
assert(stage == MESA_SHADER_TESS_CTRL);
- struct brw_tcs_prog_data *tcs_prog_data =
- (struct brw_tcs_prog_data *) prog_data;
+ struct brw_tcs_prog_data *tcs_prog_data = brw_tcs_prog_data(prog_data);
/* r1-r4 contain the ICP handles. */
payload.num_regs = 5;
}
/* Fix the disptach mask */
- if (nir->info.tcs.vertices_out % 8) {
+ if (nir->info->tcs.vertices_out % 8) {
bld.CMP(bld.null_reg_ud(), invocation_id,
- brw_imm_ud(nir->info.tcs.vertices_out), BRW_CONDITIONAL_L);
+ brw_imm_ud(nir->info->tcs.vertices_out), BRW_CONDITIONAL_L);
bld.IF(BRW_PREDICATE_NORMAL);
}
emit_nir_code();
- if (nir->info.tcs.vertices_out % 8) {
+ if (nir->info->tcs.vertices_out % 8) {
bld.emit(BRW_OPCODE_ENDIF);
}
bool
fs_visitor::run_fs(bool allow_spilling, bool do_rep_send)
{
- brw_wm_prog_data *wm_prog_data = (brw_wm_prog_data *) this->prog_data;
+ struct brw_wm_prog_data *wm_prog_data = brw_wm_prog_data(this->prog_data);
brw_wm_prog_key *wm_key = (brw_wm_prog_key *) this->key;
assert(stage == MESA_SHADER_FRAGMENT);
emit_shader_time_begin();
calculate_urb_setup();
- if (nir->info.inputs_read > 0 ||
- (nir->info.outputs_read > 0 && !wm_key->coherent_fb_fetch)) {
+ if (nir->info->inputs_read > 0 ||
+ (nir->info->outputs_read > 0 && !wm_key->coherent_fb_fetch)) {
if (devinfo->gen < 6)
emit_interpolation_setup_gen4();
else
if (devinfo->is_haswell && prog_data->total_shared > 0) {
/* Move SLM index from g0.0[27:24] to sr0.1[11:8] */
const fs_builder abld = bld.exec_all().group(1, 0);
- abld.MOV(retype(suboffset(brw_sr0_reg(), 1), BRW_REGISTER_TYPE_UW),
+ abld.MOV(retype(brw_sr0_reg(1), BRW_REGISTER_TYPE_UW),
suboffset(retype(brw_vec1_grf(0, 0), BRW_REGISTER_TYPE_UW), 1));
}
static uint8_t
computed_depth_mode(const nir_shader *shader)
{
- if (shader->info.outputs_written & BITFIELD64_BIT(FRAG_RESULT_DEPTH)) {
- switch (shader->info.fs.depth_layout) {
+ if (shader->info->outputs_written & BITFIELD64_BIT(FRAG_RESULT_DEPTH)) {
+ switch (shader->info->fs.depth_layout) {
case FRAG_DEPTH_LAYOUT_NONE:
case FRAG_DEPTH_LAYOUT_ANY:
return BRW_PSCDEPTH_ON;
}
}
-/**
- * Apply default interpolation settings to FS inputs which don't specify any.
- */
-static void
-brw_nir_set_default_interpolation(const struct gen_device_info *devinfo,
- struct nir_shader *nir,
- bool api_flat_shade,
- bool per_sample_interpolation)
-{
- assert(nir->stage == MESA_SHADER_FRAGMENT);
-
- nir_foreach_variable(var, &nir->inputs) {
- /* Apply default interpolation mode.
- *
- * Everything defaults to smooth except for the legacy GL color
- * built-in variables, which might be flat depending on API state.
- */
- if (var->data.interpolation == INTERP_MODE_NONE) {
- const bool flat = api_flat_shade &&
- (var->data.location == VARYING_SLOT_COL0 ||
- var->data.location == VARYING_SLOT_COL1);
-
- var->data.interpolation = flat ? INTERP_MODE_FLAT
- : INTERP_MODE_SMOOTH;
- }
-
- /* Apply 'sample' if necessary for API state. */
- if (per_sample_interpolation &&
- var->data.interpolation != INTERP_MODE_FLAT) {
- var->data.centroid = false;
- var->data.sample = true;
- }
-
- /* On Ironlake and below, there is only one interpolation mode.
- * Centroid interpolation doesn't mean anything on this hardware --
- * there is no multisampling.
- */
- if (devinfo->gen < 6) {
- var->data.centroid = false;
- var->data.sample = false;
- }
- }
-}
-
/**
* Demote per-sample barycentric intrinsics to centroid.
*
struct gl_program *prog,
int shader_time_index8, int shader_time_index16,
bool allow_spilling,
- bool use_rep_send,
+ bool use_rep_send, struct brw_vue_map *vue_map,
unsigned *final_assembly_size,
char **error_str)
{
nir_shader *shader = nir_shader_clone(mem_ctx, src_shader);
- shader = brw_nir_apply_sampler_key(shader, compiler->devinfo, &key->tex,
- true);
- brw_nir_set_default_interpolation(compiler->devinfo, shader,
- key->flat_shade, key->persample_interp);
- brw_nir_lower_fs_inputs(shader);
+ shader = brw_nir_apply_sampler_key(shader, compiler, &key->tex, true);
+ brw_nir_lower_fs_inputs(shader, vue_map, prog, compiler->devinfo, key);
brw_nir_lower_fs_outputs(shader);
if (!key->multisample_fbo)
NIR_PASS_V(shader, demote_sample_qualifiers);
NIR_PASS_V(shader, move_interpolation_to_top);
- shader = brw_postprocess_nir(shader, compiler->devinfo, true);
+ shader = brw_postprocess_nir(shader, 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_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);
+ 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);
+ 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);
+ (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->early_fragment_tests = shader->info.fs.early_fragment_tests;
+ 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);
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));
+ shader->info->label ?
+ shader->info->label : "unnamed",
+ shader->info->name));
}
if (simd8_cfg) {
cs_fill_push_const_info(const struct gen_device_info *devinfo,
struct brw_cs_prog_data *cs_prog_data)
{
- const struct brw_stage_prog_data *prog_data =
- (struct brw_stage_prog_data*) cs_prog_data;
+ const struct brw_stage_prog_data *prog_data = &cs_prog_data->base;
bool fill_thread_id =
cs_prog_data->thread_local_id_index >= 0 &&
cs_prog_data->thread_local_id_index < (int)prog_data->nr_params;
char **error_str)
{
nir_shader *shader = nir_shader_clone(mem_ctx, src_shader);
- shader = brw_nir_apply_sampler_key(shader, compiler->devinfo, &key->tex,
- true);
+ shader = brw_nir_apply_sampler_key(shader, compiler, &key->tex, true);
brw_nir_lower_cs_shared(shader);
prog_data->base.total_shared += shader->num_shared;
(unsigned)4 * (prog_data->thread_local_id_index + 1));
brw_nir_lower_intrinsics(shader, &prog_data->base);
- shader = brw_postprocess_nir(shader, compiler->devinfo, true);
+ shader = brw_postprocess_nir(shader, compiler, 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] = 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 =
- shader->info.cs.local_size[0] * shader->info.cs.local_size[1] *
- shader->info.cs.local_size[2];
+ 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;
unsigned simd_required = DIV_ROUND_UP(local_workgroup_size, max_cs_threads);
MESA_SHADER_COMPUTE);
if (INTEL_DEBUG & DEBUG_CS) {
char *name = ralloc_asprintf(mem_ctx, "%s compute shader %s",
- shader->info.label ? shader->info.label :
+ shader->info->label ? shader->info->label :
"unnamed",
- shader->info.name);
+ shader->info->name);
g.enable_debug(name);
}
return g.get_assembly(final_assembly_size);
}
+
+/**
+ * Test the dispatch mask packing assumptions of
+ * brw_stage_has_packed_dispatch(). Call this from e.g. the top of
+ * fs_visitor::emit_nir_code() to cause a GPU hang if any shader invocation is
+ * executed with an unexpected dispatch mask.
+ */
+static UNUSED void
+brw_fs_test_dispatch_packing(const fs_builder &bld)
+{
+ const gl_shader_stage stage = bld.shader->stage;
+
+ if (brw_stage_has_packed_dispatch(bld.shader->devinfo, stage,
+ bld.shader->stage_prog_data)) {
+ const fs_builder ubld = bld.exec_all().group(1, 0);
+ const fs_reg tmp = component(bld.vgrf(BRW_REGISTER_TYPE_UD), 0);
+ const fs_reg mask = (stage == MESA_SHADER_FRAGMENT ? brw_vmask_reg() :
+ brw_dmask_reg());
+
+ ubld.ADD(tmp, mask, brw_imm_ud(1));
+ ubld.AND(tmp, mask, tmp);
+
+ /* This will loop forever if the dispatch mask doesn't have the expected
+ * form '2^n-1', in which case tmp will be non-zero.
+ */
+ bld.emit(BRW_OPCODE_DO);
+ bld.CMP(bld.null_reg_ud(), tmp, brw_imm_ud(0), BRW_CONDITIONAL_NZ);
+ set_predicate(BRW_PREDICATE_NORMAL, bld.emit(BRW_OPCODE_WHILE));
+ }
+}
projects / mesa.git / blobdiff