fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2025-12-25 21:40:08 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions

brw: Add brw_builder::uniform()

Browse source 
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34355>
This commit is contained in:
Caio Oliveira 2025-04-03 01:14:03 -07:00 • committed by Marge Bot
parent f33d93da11
commit 7ae638c0fe
15 changed files with 71 additions and 67 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

4
src/intel/compiler/brw_builder.cpp
View file

@ -168,8 +168,8 @@ brw_emit_predicate_on_sample_mask(const brw_builder &bld, brw_inst *inst)
sample_mask.subnr == brw_flag_subreg(
subreg + inst->group / 16).subnr);
} else {
bld.group(1, 0).exec_all()
.MOV(brw_flag_subreg(subreg + inst->group / 16), sample_mask);
bld.uniform().MOV(brw_flag_subreg(subreg + inst->group / 16),
sample_mask);
}
if (inst->predicate) {

9
src/intel/compiler/brw_builder.h
View file

@ -152,6 +152,15 @@ public:
return bld;
}
/**
* Construct a builder for SIMD1 operations.
*/
brw_builder
uniform() const
{
return exec_all().group(1, 0);
}
/**
* Construct a builder for SIMD8-as-scalar
*/

7
src/intel/compiler/brw_compile_fs.cpp
View file

@ -591,7 +591,7 @@ brw_emit_repclear_shader(brw_shader &s)
for (int i = 0; i < key->nr_color_regions; ++i) {
if (i > 0)
bld.exec_all().group(1, 0).MOV(component(header, 2), brw_imm_ud(i));
bld.uniform().MOV(component(header, 2), brw_imm_ud(i));
write = bld.emit(SHADER_OPCODE_SEND);
write->resize_sources(3);
@ -1434,9 +1434,8 @@ run_fs(brw_shader &s, bool allow_spilling, bool do_rep_send)
const brw_reg dispatch_mask =
devinfo->ver >= 20 ? xe2_vec1_grf(i, 15) :
brw_vec1_grf(i + 1, 7);
bld.exec_all().group(1, 0)
.MOV(brw_sample_mask_reg(bld.group(lower_width, i)),
retype(dispatch_mask, BRW_TYPE_UW));
bld.uniform().MOV(brw_sample_mask_reg(bld.group(lower_width, i)),
retype(dispatch_mask, BRW_TYPE_UW));
}
}

2
src/intel/compiler/brw_compile_mesh.cpp
View file

@ -289,7 +289,7 @@ brw_nir_lower_mesh_primitive_count(nir_shader *nir)
static void
brw_emit_urb_fence(brw_shader &s)
{
const brw_builder bld1 = brw_builder(&s).exec_all().group(1, 0);
const brw_builder bld1 = brw_builder(&s).uniform();
brw_reg dst = bld1.vgrf(BRW_TYPE_UD);
brw_inst *fence = bld1.emit(SHADER_OPCODE_MEMORY_FENCE, dst,
brw_vec8_grf(0, 0),

17
src/intel/compiler/brw_from_nir.cpp
View file

@ -2999,7 +2999,7 @@ static void
setup_barrier_message_payload_gfx125(const brw_builder &bld,
const brw_reg &msg_payload)
{
const brw_builder ubld = bld.exec_all().group(1, 0);
const brw_builder ubld = bld.uniform();
const struct intel_device_info *devinfo = bld.shader->devinfo;
assert(devinfo->verx10 >= 125);
@ -3064,7 +3064,7 @@ emit_tcs_barrier(nir_to_brw_state &ntb)
brw_reg m0 = bld.vgrf(BRW_TYPE_UD);
brw_reg m0_2 = component(m0, 2);
const brw_builder chanbld = bld.exec_all().group(1, 0);
const brw_builder chanbld = bld.uniform();
/* Zero the message header */
bld.exec_all().MOV(m0, brw_imm_ud(0u));
@ -4462,7 +4462,7 @@ brw_from_nir_emit_fs_intrinsic(nir_to_brw_state &ntb,
const brw_reg sample_id = bld.emit_uniformize(sample_src);
const brw_reg msg_data = component(bld.group(8, 0).vgrf(BRW_TYPE_UD), 0);
bld.exec_all().group(1, 0).SHL(msg_data, sample_id, brw_imm_ud(4u));
bld.uniform().SHL(msg_data, sample_id, brw_imm_ud(4u));
brw_reg flag_reg;
struct brw_wm_prog_key *wm_prog_key = (struct brw_wm_prog_key *) s.key;
@ -4602,7 +4602,7 @@ brw_from_nir_emit_cs_intrinsic(nir_to_brw_state &ntb,
*/
if (!s.nir->info.workgroup_size_variable &&
brw_workgroup_size(s) <= s.dispatch_width) {
bld.exec_all().group(1, 0).emit(FS_OPCODE_SCHEDULING_FENCE);
bld.uniform().emit(FS_OPCODE_SCHEDULING_FENCE);
break;
}
@ -5925,7 +5925,7 @@ brw_from_nir_emit_intrinsic(nir_to_brw_state &ntb,
unsigned fence_regs_count = 0;
brw_reg fence_regs[4] = {};
const brw_builder ubld1 = bld.exec_all().group(1, 0);
const brw_builder ubld1 = bld.uniform();
/* A memory barrier with acquire semantics requires us to
* guarantee that memory operations of the specified storage
@ -7127,7 +7127,7 @@ brw_from_nir_emit_memory_access(nir_to_brw_state &ntb,
srcs[MEMORY_LOGICAL_ADDRESS] =
bld.emit_uniformize(srcs[MEMORY_LOGICAL_ADDRESS]);
const brw_builder ubld = bld.exec_all().group(1, 0);
const brw_builder ubld = bld.uniform();
unsigned total, done;
unsigned first_read_component = 0;
@ -7775,8 +7775,7 @@ emit_shader_float_controls_execution_mode(nir_to_brw_state &ntb)
if (execution_mode == FLOAT_CONTROLS_DEFAULT_FLOAT_CONTROL_MODE)
return;
brw_builder ubld = bld.exec_all().group(1, 0);
brw_builder abld = ubld.annotate("shader floats control execution mode");
brw_builder abld = bld.uniform().annotate("shader floats control execution mode");
unsigned mask, mode = brw_rnd_mode_from_nir(execution_mode, &mask);
if (mask == 0)
@ -7804,7 +7803,7 @@ brw_test_dispatch_packing(const brw_builder &bld)
if (brw_stage_has_packed_dispatch(shader->devinfo, stage,
shader->max_polygons,
shader->prog_data)) {
const brw_builder ubld = bld.exec_all().group(1, 0);
const brw_builder ubld = bld.uniform();
const brw_reg tmp = component(bld.vgrf(BRW_TYPE_UD), 0);
const brw_reg mask = uses_vmask ? brw_vmask_reg() : brw_dmask_reg();

4
src/intel/compiler/brw_lower.cpp
View file

@ -440,7 +440,7 @@ brw_lower_find_live_channel(brw_shader &s)
if (!inst->is_partial_write())
ibld.emit_undef_for_dst(inst);
const brw_builder ubld = brw_builder(inst).exec_all().group(1, 0);
const brw_builder ubld = brw_builder(inst).uniform();
brw_reg exec_mask = ubld.vgrf(BRW_TYPE_UD);
ubld.UNDEF(exec_mask);
@ -813,7 +813,7 @@ brw_lower_send_gather_inst(brw_shader &s, brw_inst *inst)
/* Fill out ARF scalar register with the physical register numbers
* and use SEND_GATHER.
*/
brw_builder ubld = brw_builder(inst).group(1, 0).exec_all();
brw_builder ubld = brw_builder(inst).uniform();
for (unsigned q = 0; q < DIV_ROUND_UP(count, 8); q++) {
uint64_t v = 0;
for (unsigned i = 0; i < 8; i++) {

12
src/intel/compiler/brw_lower_logical_sends.cpp
View file

@ -1180,7 +1180,7 @@ lower_sampler_logical_send(const brw_builder &bld, brw_inst *inst,
if (sampler_handle.file != BAD_FILE || sampler.file == IMM) {
inst->src[0] = brw_imm_ud(0);
} else {
const brw_builder ubld = bld.group(1, 0).exec_all();
const brw_builder ubld = bld.uniform();
brw_reg desc = ubld.vgrf(BRW_TYPE_UD);
ubld.SHL(desc, sampler, brw_imm_ud(8));
inst->src[0] = component(desc, 0);
@ -1199,7 +1199,7 @@ lower_sampler_logical_send(const brw_builder &bld, brw_inst *inst,
msg_type,
simd_mode,
sampler_ret_type);
const brw_builder ubld = bld.group(1, 0).exec_all();
const brw_builder ubld = bld.uniform();
brw_reg desc = ubld.vgrf(BRW_TYPE_UD);
if (surface.equals(sampler)) {
/* This case is common in GL */
@ -1334,7 +1334,7 @@ emit_predicate_on_vector_mask(const brw_builder &bld, brw_inst *inst)
bld.group() == inst->group &&
bld.dispatch_width() == inst->exec_size);
const brw_builder ubld = bld.exec_all().group(1, 0);
const brw_builder ubld = bld.uniform();
const brw_shader &s = *bld.shader;
const brw_reg vector_mask = ubld.vgrf(BRW_TYPE_UW);
@ -1386,7 +1386,7 @@ setup_surface_descriptors(const brw_builder &bld, brw_inst *inst, uint32_t desc,
inst->send_ex_bso = compiler->extended_bindless_surface_offset;
} else {
inst->desc = desc;
const brw_builder ubld = bld.exec_all().group(1, 0);
const brw_builder ubld = bld.uniform();
brw_reg tmp = ubld.vgrf(BRW_TYPE_UD);
ubld.AND(tmp, surface, brw_imm_ud(0xff));
inst->src[0] = component(tmp, 0);
@ -1424,7 +1424,7 @@ setup_lsc_surface_descriptors(const brw_builder &bld, brw_inst *inst,
if (surface.file == IMM) {
inst->src[1] = brw_imm_ud(lsc_bti_ex_desc(devinfo, surface.ud));
} else {
const brw_builder ubld = bld.exec_all().group(1, 0);
const brw_builder ubld = bld.uniform();
brw_reg tmp = ubld.vgrf(BRW_TYPE_UD);
ubld.SHL(tmp, surface, brw_imm_ud(24));
inst->src[1] = component(tmp, 0);
@ -2780,7 +2780,7 @@ brw_lower_send_descriptors(brw_shader &s)
inst->opcode != SHADER_OPCODE_SEND_GATHER)
continue;
const brw_builder ubld = brw_builder(inst).exec_all().group(1, 0);
const brw_builder ubld = brw_builder(inst).uniform();
/* Descriptor */
const unsigned rlen = inst->dst.is_null() ? 0 : inst->size_written / REG_SIZE;

9
src/intel/compiler/brw_lower_scoreboard.cpp
View file

@ -1326,8 +1326,8 @@ namespace {
/* Emit dependency into the SWSB of an extra SYNC
* instruction.
*/
const brw_builder ibld = brw_builder(inst).exec_all().group(1, 0);
brw_inst *sync = ibld.SYNC(TGL_SYNC_NOP);
const brw_builder ubld = brw_builder(inst).uniform();
brw_inst *sync = ubld.SYNC(TGL_SYNC_NOP);
sync->sched.sbid = dep.id;
sync->sched.mode = dep.unordered;
assert(!(sync->sched.mode & TGL_SBID_SET));
@ -1348,9 +1348,8 @@ namespace {
* scenario with unordered dependencies should have been
* handled above.
*/
const brw_builder ibld = brw_builder(inst)
.exec_all().group(1, 0);
brw_inst *sync = ibld.SYNC(TGL_SYNC_NOP);
const brw_builder ubld = brw_builder(inst).uniform();
brw_inst *sync = ubld.SYNC(TGL_SYNC_NOP);
sync->sched = ordered_dependency_swsb(deps[ip], jps[ip], true);
break;
}

2
src/intel/compiler/brw_lower_simd_width.cpp
View file

@ -630,7 +630,7 @@ emit_zip(const brw_builder &lbld_before, const brw_builder &lbld_after,
* have to build a single 32bit value for the SIMD32 message out of 2
* SIMD16 16 bit values.
*/
const brw_builder rbld = lbld_after.exec_all().group(1, 0);
const brw_builder rbld = lbld_after.uniform();
brw_reg local_res_reg = component(
retype(offset(tmp, lbld_before, dst_size), BRW_TYPE_UW), 0);
brw_reg final_res_reg =

4
src/intel/compiler/brw_lower_subgroup_ops.cpp
View file

@ -352,7 +352,7 @@ brw_lower_scan(brw_shader &s, brw_inst *inst)
static brw_reg
brw_fill_flag(const brw_builder &bld, unsigned v)
{
const brw_builder ubld1 = bld.exec_all().group(1, 0);
const brw_builder ubld1 = bld.uniform();
brw_reg flag = brw_flag_reg(0, 0);
if (bld.shader->dispatch_width == 32) {
@ -398,7 +398,7 @@ brw_lower_dispatch_width_vote(const brw_builder &bld, enum opcode opcode, brw_re
* TODO: Check if we still need this for newer platforms.
*/
const brw_builder ubld = devinfo->ver >= 20 ? bld.exec_all()
: bld.exec_all().group(1, 0);
: bld.uniform();
brw_reg res1 = ubld.MOV(brw_imm_d(0));
enum brw_predicate pred;

2
src/intel/compiler/brw_opt_address_reg_load.cpp
View file

@ -37,7 +37,7 @@ opt_address_reg_load_local(brw_shader &s, bblock_t *block, const brw_def_analysi
src_inst->sources > 2)
continue;
brw_builder ubld = brw_builder(&s).at(block, inst).exec_all().group(1, 0);
brw_builder ubld = brw_builder(&s).at(block, inst).uniform();
brw_reg sources[3];
for (unsigned i = 0; i < src_inst->sources; i++) {
sources[i] = inst->src[i].file == VGRF ? component(src_inst->src[i], 0) : src_inst->src[i];

24
src/intel/compiler/brw_reg_allocate.cpp
View file

@ -693,26 +693,26 @@ brw_reg_alloc::build_ex_desc(const brw_builder &bld, unsigned reg_size, bool uns
*/
brw_reg ex_desc = bld.vaddr(BRW_TYPE_UD,
BRW_ADDRESS_SUBREG_INDIRECT_SPILL_DESC);
brw_inst *inst = bld.exec_all().group(1, 0).AND(
ex_desc,
retype(brw_vec1_grf(0, 5), BRW_TYPE_UD),
brw_imm_ud(INTEL_MASK(31, 10)));
brw_builder ubld = bld.uniform();
brw_inst *inst = ubld.AND(ex_desc,
retype(brw_vec1_grf(0, 5), BRW_TYPE_UD),
brw_imm_ud(INTEL_MASK(31, 10)));
_mesa_set_add(spill_insts, inst);
const intel_device_info *devinfo = bld.shader->devinfo;
if (devinfo->verx10 >= 200) {
inst = bld.exec_all().group(1, 0).SHR(
ex_desc, ex_desc, brw_imm_ud(4));
inst = ubld.SHR(ex_desc, ex_desc, brw_imm_ud(4));
_mesa_set_add(spill_insts, inst);
} else {
if (unspill) {
inst = bld.exec_all().group(1, 0).OR(
ex_desc, ex_desc, brw_imm_ud(BRW_SFID_UGM));
inst = ubld.OR(ex_desc, ex_desc, brw_imm_ud(BRW_SFID_UGM));
_mesa_set_add(spill_insts, inst);
} else {
inst = bld.exec_all().group(1, 0).OR(
ex_desc, ex_desc,
brw_imm_ud(brw_message_ex_desc(devinfo, reg_size) | BRW_SFID_UGM));
inst = ubld.OR(ex_desc,
ex_desc,
brw_imm_ud(brw_message_ex_desc(devinfo, reg_size) | BRW_SFID_UGM));
_mesa_set_add(spill_insts, inst);
}
}
@ -816,7 +816,7 @@ brw_reg_alloc::emit_unspill(const brw_builder &bld,
const bool use_transpose =
bld.dispatch_width() > 16 * reg_unit(devinfo) ||
bld.has_writemask_all();
const brw_builder ubld = use_transpose ? bld.exec_all().group(1, 0) : bld;
const brw_builder ubld = use_transpose ? bld.uniform() : bld;
brw_reg offset;
if (use_transpose) {
offset = build_single_offset(ubld, spill_offset, ip);

6
src/intel/compiler/brw_workaround.cpp
View file

@ -98,8 +98,7 @@ brw_workaround_memory_fence_before_eot(brw_shader &s)
if (!has_ugm_write_or_atomic)
break;
const brw_builder ibld(inst);
const brw_builder ubld = ibld.exec_all().group(1, 0);
const brw_builder ubld = brw_builder(inst).uniform();
brw_reg dst = ubld.vgrf(BRW_TYPE_UD);
brw_inst *dummy_fence = ubld.emit(SHADER_OPCODE_MEMORY_FENCE,
@ -342,8 +341,7 @@ brw_workaround_source_arf_before_eot(brw_shader &s)
*/
assert(++eot_count == 1);
const brw_builder ibld(inst);
const brw_builder ubld = ibld.exec_all().group(1, 0);
const brw_builder ubld = brw_builder(inst).uniform();
if (flags_unread & 0x0f)
ubld.MOV(ubld.null_reg_ud(), retype(brw_flag_reg(0, 0), BRW_TYPE_UD));

30
src/intel/compiler/test_lower_scoreboard.cpp
View file

@ -35,7 +35,7 @@ protected:
brw_inst *
SYNC_NOP(const brw_builder &bld)
{
return bld.group(1, 0).exec_all().SYNC(TGL_SYNC_NOP);
return bld.uniform().SYNC(TGL_SYNC_NOP);
}
brw_inst *
@ -876,13 +876,13 @@ TEST_F(scoreboard_test, gitlab_issue_from_mr_29723)
brw_reg a = brw_ud8_grf(29, 0);
brw_reg b = brw_ud8_grf(2, 0);
auto bld1 = bld.exec_all().group(1, 0);
auto bld1 = bld.uniform();
bld1.ADD( a, stride(b, 0, 1, 0), brw_imm_ud(256));
bld1.CMP(brw_null_reg(), stride(a, 2, 1, 2), stride(b, 0, 1, 0), BRW_CONDITIONAL_L);
EXPECT_PROGRESS(brw_lower_scoreboard, bld);
auto exp1 = exp.exec_all().group(1, 0);
auto exp1 = exp.uniform();
exp1.ADD( a, stride(b, 0, 1, 0), brw_imm_ud(256));
exp1.CMP(brw_null_reg(), stride(a, 2, 1, 2), stride(b, 0, 1, 0), BRW_CONDITIONAL_L)->sched = SWSB("@1");
@ -969,13 +969,13 @@ TEST_F(scoreboard_test, gitlab_issue_11069)
brw_reg a = brw_ud8_grf(76, 0);
brw_reg b = brw_ud8_grf(2, 0);
auto bld1 = bld.exec_all().group(1, 0);
auto bld1 = bld.uniform();
bld1.ADD(stride(a, 2, 1, 2), stride(b, 0, 1, 0), brw_imm_ud(0x80));
bld1.CMP( brw_null_reg(), stride(a, 0, 1, 0), stride(b, 0, 1, 0), BRW_CONDITIONAL_L);
EXPECT_PROGRESS(brw_lower_scoreboard, bld);
auto exp1 = exp.exec_all().group(1, 0);
auto exp1 = exp.uniform();
exp1.ADD(stride(a, 2, 1, 2), stride(b, 0, 1, 0), brw_imm_ud(0x80));
exp1.CMP( brw_null_reg(), stride(a, 0, 1, 0), stride(b, 0, 1, 0), BRW_CONDITIONAL_L)->sched = SWSB("@1");
@ -1109,14 +1109,14 @@ TEST_F(scoreboard_test, scalar_register_mov_immediate_is_in_scalar_pipe)
brw_reg imm = brw_imm_uw(0x1415);
brw_reg r20 = brw_uw8_grf(20, 0);
bld.group(1, 0).exec_all().MOV(scalar, imm);
bld .MOV(r20, scalar);
bld.uniform().MOV(scalar, imm);
bld .MOV(r20, scalar);
EXPECT_PROGRESS(brw_lower_scoreboard, bld);
exp.group(1, 0).exec_all().MOV(scalar, imm);
SYNC_NOP(exp )->sched = SWSB("S@1");
exp .MOV(r20, scalar);
exp.uniform().MOV(scalar, imm);
SYNC_NOP(exp )->sched = SWSB("S@1");
exp .MOV(r20, scalar);
EXPECT_SHADERS_MATCH(bld, exp);
}
@ -1132,14 +1132,14 @@ TEST_F(scoreboard_test, scalar_register_mov_grf_is_not_in_scalar_pipe)
brw_reg r10 = brw_uw8_grf(10, 0);
brw_reg r20 = brw_uw8_grf(20, 0);
bld.group(1, 0).exec_all().MOV(scalar, r10);
bld .MOV(r20, scalar);
bld.uniform().MOV(scalar, r10);
bld .MOV(r20, scalar);
EXPECT_PROGRESS(brw_lower_scoreboard, bld);
exp.group(1, 0).exec_all().MOV (scalar, r10);
SYNC_NOP(exp )->sched = SWSB("I@1");
exp .MOV (r20, scalar);
exp.uniform().MOV (scalar, r10);
SYNC_NOP(exp )->sched = SWSB("I@1");
exp .MOV (r20, scalar);
EXPECT_SHADERS_MATCH(bld, exp);
}

6
src/intel/compiler/test_opt_combine_constants.cpp
View file

@ -23,8 +23,8 @@ TEST_F(CombineConstantsTest, Simple)
brw_reg tmp = component(exp.vgrf(BRW_TYPE_D), 0);
exp.group(1, 0).exec_all().MOV(tmp, imm_a);
exp .SEL(r, tmp, imm_b);
exp.uniform().MOV(tmp, imm_a);
exp .SEL(r, tmp, imm_b);
EXPECT_SHADERS_MATCH(bld, exp);
}
@ -51,7 +51,7 @@ TEST_F(CombineConstantsTest, DoContainingDo)
/* Explicit emit the expected FLOW instruction. */
exp.emit(BRW_OPCODE_DO);
brw_reg tmp = component(exp.vgrf(BRW_TYPE_D), 0);
exp.group(1, 0).exec_all().MOV(tmp, imm_a);
exp.uniform().MOV(tmp, imm_a);
exp.emit(SHADER_OPCODE_FLOW);
exp.DO();
exp.SEL(r1, tmp, imm_b);