X-Git-Url: https://git.libre-soc.org/?a=blobdiff_plain;f=src%2Fmesa%2Fdrivers%2Fdri%2Fi965%2Fbrw_fs.cpp;h=68e73cc5cd801a672bbb73e884444c5f9e6556d6;hb=a497ab6838ae5a9898abfed82f7bc8295b490911;hp=0244f593149bac3d7e7745a107f06e47a7277fee;hpb=69570bbad876bb9da609c3b651aacda28cecc542;p=mesa.git diff --git a/src/mesa/drivers/dri/i965/brw_fs.cpp b/src/mesa/drivers/dri/i965/brw_fs.cpp index 0244f593149..68e73cc5cd8 100644 --- a/src/mesa/drivers/dri/i965/brw_fs.cpp +++ b/src/mesa/drivers/dri/i965/brw_fs.cpp @@ -172,12 +172,12 @@ fs_visitor::VARYING_PULL_CONSTANT_LOAD(const fs_builder &bld, * 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)); @@ -191,7 +191,7 @@ fs_visitor::VARYING_PULL_CONSTANT_LOAD(const fs_builder &bld, 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( @@ -240,12 +240,6 @@ fs_inst::equals(fs_inst *inst) const 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 { @@ -353,7 +347,7 @@ fs_inst::is_copy_payload(const brw::simple_allocator &grf_alloc) 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) @@ -441,15 +435,6 @@ fs_reg::equals(const fs_reg &r) const 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 { @@ -477,6 +462,8 @@ type_size_scalar(const struct glsl_type *type) case GLSL_TYPE_BOOL: return type->components(); case GLSL_TYPE_DOUBLE: + case GLSL_TYPE_UINT64: + case GLSL_TYPE_INT64: return type->components() * 2; case GLSL_TYPE_ARRAY: return type_size_scalar(type->fields.array) * type->length; @@ -507,32 +494,6 @@ type_size_scalar(const struct glsl_type *type) 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. * @@ -562,15 +523,14 @@ fs_visitor::get_timestamp(const fs_builder &bld) 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 @@ -581,8 +541,7 @@ fs_visitor::emit_shader_time_end() 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()). @@ -591,22 +550,21 @@ fs_visitor::emit_shader_time_end() * 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); @@ -761,7 +719,7 @@ fs_inst::components_read(unsigned i) const 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) @@ -817,8 +775,8 @@ fs_inst::components_read(unsigned i) const } } -int -fs_inst::regs_read(int arg) const +unsigned +fs_inst::size_read(int arg) const { switch (opcode) { case FS_OPCODE_FB_WRITE: @@ -837,78 +795,52 @@ fs_inst::regs_read(int arg) const case SHADER_OPCODE_TYPED_SURFACE_WRITE: case FS_OPCODE_INTERPOLATE_AT_PER_SLOT_OFFSET: if (arg == 0) - return mlen; + return mlen * REG_SIZE; break; case FS_OPCODE_UNIFORM_PULL_CONSTANT_LOAD_GEN7: /* The payload is actually stored in src1 */ if (arg == 1) - return mlen; + return mlen * REG_SIZE; break; case FS_OPCODE_LINTERP: if (arg == 1) - return 1; + return 16; break; case SHADER_OPCODE_LOAD_PAYLOAD: if (arg < this->header_size) - return 1; + return REG_SIZE; break; case CS_OPCODE_CS_TERMINATE: case SHADER_OPCODE_BARRIER: - return 1; + return REG_SIZE; 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 / 4; - } 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 DIV_ROUND_UP(region_length, REG_SIZE); - } else { - assert(!"Invalid register file"); - } + return src[2].ud; } break; default: if (is_tex() && arg == 0 && src[0].file == VGRF) - return mlen; + return mlen * REG_SIZE; break; } switch (src[arg].file) { case UNIFORM: case IMM: - return 1; + return components_read(arg) * type_sz(src[arg].type); case BAD_FILE: case ARF: case FIXED_GRF: case VGRF: case ATTR: - return DIV_ROUND_UP(components_read(arg) * - src[arg].component_size(exec_size), - REG_SIZE); + return components_read(arg) * src[arg].component_size(exec_size); case MRF: unreachable("MRF registers are not allowed as sources"); } @@ -1169,7 +1101,7 @@ void 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) { @@ -1289,9 +1221,9 @@ fs_visitor::emit_sampleid_setup() 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 @@ -1337,7 +1269,7 @@ fs_reg * 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)); @@ -1388,7 +1320,7 @@ fs_visitor::resolve_source_modifiers(const fs_reg &src) 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. @@ -1405,8 +1337,7 @@ fs_visitor::emit_gs_thread_end() { 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); @@ -1495,7 +1426,7 @@ void 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, @@ -1504,7 +1435,7 @@ fs_visitor::calculate_urb_setup() 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. @@ -1516,14 +1447,14 @@ fs_visitor::calculate_urb_setup() * 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 @@ -1533,7 +1464,7 @@ fs_visitor::calculate_urb_setup() 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; @@ -1542,7 +1473,7 @@ fs_visitor::calculate_urb_setup() 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; } @@ -1575,7 +1506,7 @@ fs_visitor::calculate_urb_setup() * * 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++; } @@ -1586,7 +1517,7 @@ void 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; @@ -1653,7 +1584,7 @@ fs_visitor::convert_attr_sources_to_hw_regs(fs_inst *inst) 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); @@ -1684,7 +1615,7 @@ fs_visitor::assign_tes_urb_setup() { 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; @@ -1699,10 +1630,10 @@ fs_visitor::assign_gs_urb_setup() { 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. */ @@ -1732,6 +1663,12 @@ fs_visitor::assign_gs_urb_setup() 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 */ @@ -1916,7 +1853,10 @@ fs_visitor::compact_virtual_grfs() } static void -set_push_pull_constant_loc(unsigned uniform, int *chunk_start, bool contiguous, +set_push_pull_constant_loc(unsigned uniform, int *chunk_start, + unsigned *max_chunk_bitsize, + bool contiguous, unsigned bitsize, + const unsigned target_bitsize, int *push_constant_loc, int *pull_constant_loc, unsigned *num_push_constants, unsigned *num_pull_constants, @@ -1928,11 +1868,23 @@ set_push_pull_constant_loc(unsigned uniform, int *chunk_start, bool contiguous, if (*chunk_start < 0) *chunk_start = uniform; + /* Keep track of the maximum bit size access in contiguous uniforms */ + *max_chunk_bitsize = MAX2(*max_chunk_bitsize, bitsize); + /* If this element does not need to be contiguous with the next, we * split at this point and everything between chunk_start and u forms a * single chunk. */ if (!contiguous) { + /* If bitsize doesn't match the target one, skip it */ + if (*max_chunk_bitsize != target_bitsize) { + /* FIXME: right now we only support 32 and 64-bit accesses */ + assert(*max_chunk_bitsize == 4 || *max_chunk_bitsize == 8); + *max_chunk_bitsize = 0; + *chunk_start = -1; + return; + } + unsigned chunk_size = uniform - *chunk_start + 1; /* Decide whether we should push or pull this parameter. In the @@ -1950,6 +1902,7 @@ set_push_pull_constant_loc(unsigned uniform, int *chunk_start, bool contiguous, pull_constant_loc[j] = (*num_pull_constants)++; } + *max_chunk_bitsize = 0; *chunk_start = -1; } } @@ -1972,8 +1925,8 @@ fs_visitor::assign_constant_locations() bool is_live[uniforms]; memset(is_live, 0, sizeof(is_live)); - bool is_live_64bit[uniforms]; - memset(is_live_64bit, 0, sizeof(is_live_64bit)); + unsigned bitsize_access[uniforms]; + memset(bitsize_access, 0, sizeof(bitsize_access)); /* For each uniform slot, a value of true indicates that the given slot and * the next slot must remain contiguous. This is used to keep us from @@ -1984,7 +1937,7 @@ fs_visitor::assign_constant_locations() 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: * @@ -2010,20 +1963,18 @@ fs_visitor::assign_constant_locations() for (unsigned j = constant_nr; j < last; j++) { is_live[j] = true; contiguous[j] = true; - if (type_sz(inst->src[i].type) == 8) { - is_live_64bit[j] = true; - } + bitsize_access[j] = MAX2(bitsize_access[j], type_sz(inst->src[i].type)); } is_live[last] = true; + bitsize_access[last] = MAX2(bitsize_access[last], type_sz(inst->src[i].type)); } else { if (constant_nr >= 0 && constant_nr < (int) uniforms) { int regs_read = inst->components_read(i) * type_sz(inst->src[i].type) / 4; for (int j = 0; j < regs_read; j++) { is_live[constant_nr + j] = true; - if (type_sz(inst->src[i].type) == 8) { - is_live_64bit[constant_nr + j] = true; - } + bitsize_access[constant_nr + j] = + MAX2(bitsize_access[constant_nr + j], type_sz(inst->src[i].type)); } } } @@ -2062,13 +2013,17 @@ fs_visitor::assign_constant_locations() memset(pull_constant_loc, -1, uniforms * sizeof(*pull_constant_loc)); int chunk_start = -1; + unsigned max_chunk_bitsize = 0; /* First push 64-bit uniforms to ensure they are properly aligned */ + const unsigned uniform_64_bit_size = type_sz(BRW_REGISTER_TYPE_DF); for (unsigned u = 0; u < uniforms; u++) { - if (!is_live[u] || !is_live_64bit[u]) + if (!is_live[u]) continue; - set_push_pull_constant_loc(u, &chunk_start, contiguous[u], + set_push_pull_constant_loc(u, &chunk_start, &max_chunk_bitsize, + contiguous[u], bitsize_access[u], + uniform_64_bit_size, push_constant_loc, pull_constant_loc, &num_push_constants, &num_pull_constants, max_push_components, max_chunk_size, @@ -2077,15 +2032,18 @@ fs_visitor::assign_constant_locations() } /* Then push the rest of uniforms */ + const unsigned uniform_32_bit_size = type_sz(BRW_REGISTER_TYPE_F); for (unsigned u = 0; u < uniforms; u++) { - if (!is_live[u] || is_live_64bit[u]) + if (!is_live[u]) continue; /* Skip thread_local_id_index to put it in the last push register. */ if (thread_local_id_index == (int)u) continue; - set_push_pull_constant_loc(u, &chunk_start, contiguous[u], + set_push_pull_constant_loc(u, &chunk_start, &max_chunk_bitsize, + contiguous[u], bitsize_access[u], + uniform_32_bit_size, push_constant_loc, pull_constant_loc, &num_push_constants, &num_pull_constants, max_push_components, max_chunk_size, @@ -2106,7 +2064,7 @@ fs_visitor::assign_constant_locations() 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. * @@ -2129,7 +2087,7 @@ fs_visitor::assign_constant_locations() 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; } @@ -2163,27 +2121,22 @@ fs_visitor::lower_constant_loads() 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); } @@ -2547,7 +2500,7 @@ fs_visitor::opt_sampler_eot() 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->regs_read(i) * REG_SIZE != tex_inst->size_written) + 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) @@ -2678,16 +2631,18 @@ fs_visitor::opt_redundant_discard_jumps() /** * 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 @@ -2730,7 +2685,7 @@ fs_visitor::compute_to_mrf() foreach_inst_in_block_reverse_starting_from(fs_inst, scan_inst, inst) { if (regions_overlap(scan_inst->dst, scan_inst->size_written, - inst->src[0], inst->regs_read(0) * REG_SIZE)) { + inst->src[0], inst->size_read(0))) { /* Found the last thing to write our reg we want to turn * into a compute-to-MRF. */ @@ -2747,9 +2702,8 @@ fs_visitor::compute_to_mrf() * 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 + inst->regs_read(0)) + 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. */ @@ -2767,8 +2721,7 @@ fs_visitor::compute_to_mrf() /* 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; } @@ -2785,8 +2738,8 @@ fs_visitor::compute_to_mrf() */ bool interfered = false; for (int i = 0; i < scan_inst->sources; i++) { - if (regions_overlap(scan_inst->src[i], scan_inst->regs_read(i) * REG_SIZE, - inst->src[0], inst->regs_read(0) * REG_SIZE)) { + if (regions_overlap(scan_inst->src[i], scan_inst->size_read(i), + inst->src[0], inst->size_read(0))) { interfered = true; } } @@ -2823,21 +2776,20 @@ fs_visitor::compute_to_mrf() foreach_inst_in_block_reverse_starting_from(fs_inst, scan_inst, inst) { if (regions_overlap(scan_inst->dst, scan_inst->size_written, - inst->src[0], inst->regs_read(0) * REG_SIZE)) { + 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. @@ -2849,11 +2801,11 @@ fs_visitor::compute_to_mrf() /* 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; @@ -2882,6 +2834,14 @@ fs_visitor::eliminate_find_live_channel() 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: @@ -3027,7 +2987,7 @@ fs_visitor::remove_duplicate_mrf_writes() if (last_mrf_move[i] && regions_overlap(inst->dst, inst->size_written, last_mrf_move[i]->src[0], - last_mrf_move[i]->regs_read(0) * REG_SIZE)) { + last_mrf_move[i]->size_read(0))) { last_mrf_move[i] = NULL; } } @@ -3213,10 +3173,6 @@ fs_visitor::insert_gen4_send_dependency_workarounds() 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); @@ -3253,44 +3209,18 @@ fs_visitor::lower_uniform_pull_constant_loads() 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; - } + const fs_builder ubld = fs_builder(this, block, inst).exec_all(); + const fs_reg payload = ubld.group(8, 0).vgrf(BRW_REGISTER_TYPE_UD); - /* 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; - - 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 { @@ -3500,62 +3430,27 @@ fs_visitor::lower_integer_multiplication() 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, @@ -3675,7 +3570,7 @@ setup_color_payload(const fs_builder &bld, const brw_wm_prog_key *key, 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) { @@ -3759,6 +3654,12 @@ lower_fb_write_logical_send(const fs_builder &bld, fs_inst *inst, */ 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); @@ -3913,8 +3814,8 @@ lower_sampler_logical_send_gen4(const fs_builder &bld, fs_inst *inst, opcode op, } 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); @@ -3957,7 +3858,6 @@ lower_sampler_logical_send_gen5(const fs_builder &bld, fs_inst *inst, opcode op, const fs_reg &sample_index, const fs_reg &surface, const fs_reg &sampler, - const fs_reg &offset_value, unsigned coord_components, unsigned grad_components) { @@ -3965,7 +3865,7 @@ lower_sampler_logical_send_gen5(const fs_builder &bld, fs_inst *inst, opcode op, 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. */ @@ -4065,7 +3965,7 @@ lower_sampler_logical_send_gen7(const fs_builder &bld, fs_inst *inst, opcode op, 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) { @@ -4077,7 +3977,7 @@ lower_sampler_logical_send_gen7(const fs_builder &bld, fs_inst *inst, opcode op, 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 @@ -4222,7 +4122,7 @@ lower_sampler_logical_send_gen7(const fs_builder &bld, fs_inst *inst, opcode op, 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)); @@ -4274,7 +4174,7 @@ lower_sampler_logical_send(const fs_builder &bld, fs_inst *inst, opcode op) 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); @@ -4283,12 +4183,12 @@ lower_sampler_logical_send(const fs_builder &bld, fs_inst *inst, opcode op) 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, @@ -4431,7 +4331,7 @@ fs_visitor::lower_logical_sends() 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; @@ -4607,7 +4507,7 @@ get_fpu_lowered_simd_width(const struct gen_device_info *devinfo, unsigned reg_count = DIV_ROUND_UP(inst->size_written, REG_SIZE); for (unsigned i = 0; i < inst->sources; i++) - reg_count = MAX2(reg_count, (unsigned)inst->regs_read(i)); + reg_count = MAX2(reg_count, DIV_ROUND_UP(inst->size_read(i), REG_SIZE)); /* Calculate the maximum execution size of the instruction based on the * factor by which it goes over the hardware limit of 2 GRFs. @@ -4631,8 +4531,8 @@ get_fpu_lowered_simd_width(const struct gen_device_info *devinfo, */ if (devinfo->gen < 8) { for (unsigned i = 0; i < inst->sources; i++) { - if (DIV_ROUND_UP(inst->size_written, REG_SIZE) == 2 && - inst->regs_read(i) != 0 && inst->regs_read(i) != 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)) { @@ -4757,7 +4657,7 @@ get_sampler_lowered_simd_width(const struct gen_device_info *devinfo, 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 @@ -5114,7 +5014,7 @@ needs_dst_copy(const fs_builder &lbld, const fs_inst *inst) * the data read from the same source by other lowered instructions. */ if (regions_overlap(inst->dst, inst->size_written, - inst->src[i], inst->regs_read(i) * REG_SIZE) && + inst->src[i], inst->size_read(i)) && !inst->dst.equals(inst->src[i])) return true; } @@ -5315,10 +5215,6 @@ fs_visitor::dump_instruction(backend_instruction *be_inst, FILE *file) 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); @@ -5330,10 +5226,10 @@ fs_visitor::dump_instruction(backend_instruction *be_inst, FILE *file) 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) { @@ -5353,12 +5249,19 @@ fs_visitor::dump_instruction(backend_instruction *be_inst, FILE *file) 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)); @@ -5371,10 +5274,6 @@ fs_visitor::dump_instruction(backend_instruction *be_inst, FILE *file) switch (inst->src[i].file) { case VGRF: fprintf(file, "vgrf%d", inst->src[i].nr); - if (alloc.sizes[inst->src[i].nr] != (unsigned)inst->regs_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); @@ -5383,14 +5282,10 @@ fs_visitor::dump_instruction(backend_instruction *be_inst, FILE *file) 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)"); @@ -5441,10 +5336,17 @@ fs_visitor::dump_instruction(backend_instruction *be_inst, FILE *file) 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, "|"); @@ -5508,11 +5410,7 @@ void 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); @@ -5528,7 +5426,7 @@ fs_visitor::setup_fs_payload_gen6() * 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) { @@ -5539,7 +5437,7 @@ fs_visitor::setup_fs_payload_gen6() /* 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++; @@ -5551,7 +5449,7 @@ fs_visitor::setup_fs_payload_gen6() /* 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++; @@ -5563,7 +5461,7 @@ fs_visitor::setup_fs_payload_gen6() /* 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 @@ -5580,7 +5478,7 @@ fs_visitor::setup_fs_payload_gen6() /* 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; @@ -5594,7 +5492,7 @@ fs_visitor::setup_fs_payload_gen6() /* 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; } } @@ -5611,10 +5509,8 @@ fs_visitor::setup_gs_payload() { 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; @@ -5633,15 +5529,15 @@ fs_visitor::setup_gs_payload() * Note that the GS reads 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; } } @@ -5742,7 +5638,7 @@ fs_visitor::optimize() 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); \ } \ @@ -5756,7 +5652,7 @@ fs_visitor::optimize() 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); } @@ -5776,7 +5672,7 @@ fs_visitor::optimize() 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); @@ -5800,7 +5696,7 @@ fs_visitor::optimize() } if (OPT(lower_d2x)) { - OPT(opt_copy_propagate); + OPT(opt_copy_propagation); OPT(dead_code_eliminate); } @@ -5812,12 +5708,12 @@ fs_visitor::optimize() 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 @@ -5846,7 +5742,7 @@ fs_visitor::optimize() 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); } @@ -5946,7 +5842,7 @@ fs_visitor::allocate_registers(bool allow_spilling) 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); @@ -6022,8 +5918,7 @@ fs_visitor::run_tcs_single_patch() { 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; @@ -6053,15 +5948,15 @@ fs_visitor::run_tcs_single_patch() } /* Fix the disptach mask */ - if (nir->info.tcs.vertices_out % 8) { + if (nir->info->tess.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->tess.tcs_vertices_out), BRW_CONDITIONAL_L); bld.IF(BRW_PREDICATE_NORMAL); } emit_nir_code(); - if (nir->info.tcs.vertices_out % 8) { + if (nir->info->tess.tcs_vertices_out % 8) { bld.emit(BRW_OPCODE_ENDIF); } @@ -6184,7 +6079,7 @@ fs_visitor::run_gs() 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); @@ -6204,8 +6099,8 @@ fs_visitor::run_fs(bool allow_spilling, bool do_rep_send) 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 @@ -6269,7 +6164,7 @@ fs_visitor::run_cs() 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)); } @@ -6369,8 +6264,8 @@ brw_compute_flat_inputs(struct brw_wm_prog_data *prog_data, 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; @@ -6456,50 +6351,6 @@ move_interpolation_to_top(nir_shader *nir) } } -/** - * 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. * @@ -6550,41 +6401,48 @@ brw_compile_fs(const struct brw_compiler *compiler, void *log_data, 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) { + const struct gen_device_info *devinfo = compiler->devinfo; + 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, devinfo, key); brw_nir_lower_fs_outputs(shader); + + if (devinfo->gen < 6) { + brw_setup_vue_interpolation(vue_map, shader, prog_data, devinfo); + } + 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); @@ -6667,9 +6525,9 @@ brw_compile_fs(const struct brw_compiler *compiler, void *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)); + shader->info->label ? + shader->info->label : "unnamed", + shader->info->name)); } if (simd8_cfg) { @@ -6724,8 +6582,7 @@ static void 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; @@ -6787,8 +6644,7 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data, 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; @@ -6801,14 +6657,14 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data, (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); @@ -6898,9 +6754,9 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data, MESA_SHADER_COMPUTE); if (INTEL_DEBUG & DEBUG_CS) { char *name = ralloc_asprintf(mem_ctx, "%s compute shader %s", - shader->info.label ? shader->info.label : + shader->info->label ? shader->info->label : "unnamed", - shader->info.name); + shader->info->name); g.enable_debug(name); } @@ -6908,3 +6764,33 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data, 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)); + } +}