From 7ae638c0fe049681cd487367e7f48a9e2ff5112d Mon Sep 17 00:00:00 2001 From: Caio Oliveira Date: Thu, 3 Apr 2025 01:14:03 -0700 Subject: [PATCH] brw: Add brw_builder::uniform() Reviewed-by: Lionel Landwerlin Reviewed-by: Kenneth Graunke Part-of: --- src/intel/compiler/brw_builder.cpp | 4 +-- src/intel/compiler/brw_builder.h | 9 ++++++ src/intel/compiler/brw_compile_fs.cpp | 7 ++--- src/intel/compiler/brw_compile_mesh.cpp | 2 +- src/intel/compiler/brw_from_nir.cpp | 17 +++++------ src/intel/compiler/brw_lower.cpp | 4 +-- .../compiler/brw_lower_logical_sends.cpp | 12 ++++---- src/intel/compiler/brw_lower_scoreboard.cpp | 9 +++--- src/intel/compiler/brw_lower_simd_width.cpp | 2 +- src/intel/compiler/brw_lower_subgroup_ops.cpp | 4 +-- .../compiler/brw_opt_address_reg_load.cpp | 2 +- src/intel/compiler/brw_reg_allocate.cpp | 24 +++++++-------- src/intel/compiler/brw_workaround.cpp | 6 ++-- src/intel/compiler/test_lower_scoreboard.cpp | 30 +++++++++---------- .../compiler/test_opt_combine_constants.cpp | 6 ++-- 15 files changed, 71 insertions(+), 67 deletions(-) diff --git a/src/intel/compiler/brw_builder.cpp b/src/intel/compiler/brw_builder.cpp index 5643bbacc95..785616f8c59 100644 --- a/src/intel/compiler/brw_builder.cpp +++ b/src/intel/compiler/brw_builder.cpp @@ -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) { diff --git a/src/intel/compiler/brw_builder.h b/src/intel/compiler/brw_builder.h index b191b6c5735..8313c4ee185 100644 --- a/src/intel/compiler/brw_builder.h +++ b/src/intel/compiler/brw_builder.h @@ -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 */ diff --git a/src/intel/compiler/brw_compile_fs.cpp b/src/intel/compiler/brw_compile_fs.cpp index 884a7865664..1743184b4e9 100644 --- a/src/intel/compiler/brw_compile_fs.cpp +++ b/src/intel/compiler/brw_compile_fs.cpp @@ -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)); } } diff --git a/src/intel/compiler/brw_compile_mesh.cpp b/src/intel/compiler/brw_compile_mesh.cpp index 4c65d569a5e..fae56100f2a 100644 --- a/src/intel/compiler/brw_compile_mesh.cpp +++ b/src/intel/compiler/brw_compile_mesh.cpp @@ -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), diff --git a/src/intel/compiler/brw_from_nir.cpp b/src/intel/compiler/brw_from_nir.cpp index fedd4a65688..414440bab9d 100644 --- a/src/intel/compiler/brw_from_nir.cpp +++ b/src/intel/compiler/brw_from_nir.cpp @@ -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(); diff --git a/src/intel/compiler/brw_lower.cpp b/src/intel/compiler/brw_lower.cpp index ee29dc14b10..baa126c65b0 100644 --- a/src/intel/compiler/brw_lower.cpp +++ b/src/intel/compiler/brw_lower.cpp @@ -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++) { diff --git a/src/intel/compiler/brw_lower_logical_sends.cpp b/src/intel/compiler/brw_lower_logical_sends.cpp index 455f6f78066..466fe3b4bc7 100644 --- a/src/intel/compiler/brw_lower_logical_sends.cpp +++ b/src/intel/compiler/brw_lower_logical_sends.cpp @@ -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; diff --git a/src/intel/compiler/brw_lower_scoreboard.cpp b/src/intel/compiler/brw_lower_scoreboard.cpp index 730477177c7..a3a08abaa00 100644 --- a/src/intel/compiler/brw_lower_scoreboard.cpp +++ b/src/intel/compiler/brw_lower_scoreboard.cpp @@ -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; } diff --git a/src/intel/compiler/brw_lower_simd_width.cpp b/src/intel/compiler/brw_lower_simd_width.cpp index 2fbd02d767f..2404ad09453 100644 --- a/src/intel/compiler/brw_lower_simd_width.cpp +++ b/src/intel/compiler/brw_lower_simd_width.cpp @@ -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 = diff --git a/src/intel/compiler/brw_lower_subgroup_ops.cpp b/src/intel/compiler/brw_lower_subgroup_ops.cpp index e0496bab3c2..45d00726413 100644 --- a/src/intel/compiler/brw_lower_subgroup_ops.cpp +++ b/src/intel/compiler/brw_lower_subgroup_ops.cpp @@ -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; diff --git a/src/intel/compiler/brw_opt_address_reg_load.cpp b/src/intel/compiler/brw_opt_address_reg_load.cpp index 90d246c8d04..f08822d5695 100644 --- a/src/intel/compiler/brw_opt_address_reg_load.cpp +++ b/src/intel/compiler/brw_opt_address_reg_load.cpp @@ -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]; diff --git a/src/intel/compiler/brw_reg_allocate.cpp b/src/intel/compiler/brw_reg_allocate.cpp index 61ff23a34d1..db3db369c24 100644 --- a/src/intel/compiler/brw_reg_allocate.cpp +++ b/src/intel/compiler/brw_reg_allocate.cpp @@ -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); diff --git a/src/intel/compiler/brw_workaround.cpp b/src/intel/compiler/brw_workaround.cpp index c3ef97c5026..647dd3f19f7 100644 --- a/src/intel/compiler/brw_workaround.cpp +++ b/src/intel/compiler/brw_workaround.cpp @@ -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)); diff --git a/src/intel/compiler/test_lower_scoreboard.cpp b/src/intel/compiler/test_lower_scoreboard.cpp index eb05fd84d95..cb938166163 100644 --- a/src/intel/compiler/test_lower_scoreboard.cpp +++ b/src/intel/compiler/test_lower_scoreboard.cpp @@ -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); } diff --git a/src/intel/compiler/test_opt_combine_constants.cpp b/src/intel/compiler/test_opt_combine_constants.cpp index 885ceff11fd..75f0a323237 100644 --- a/src/intel/compiler/test_opt_combine_constants.cpp +++ b/src/intel/compiler/test_opt_combine_constants.cpp @@ -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);