OpenHarmony/third_party_mesa3d
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

nir: Rename nir_intrinsic_barrier to control_barrier

Browse Source 
This is a more explicit name now that we don't want it to be doing any
memory barrier stuff for us.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3307>
This commit is contained in:
Jason Ekstrand
2020-01-07 14:54:26 -06:00
committed by Marge Bot
parent bd3ab75aef
commit e40b11bbcb
18 changed files with 30 additions and 24 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
src/amd/compiler/aco_instruction_selection.cpp
View File

@@ -5714,7 +5714,7 @@ void visit_intrinsic(isel_context *ctx, nir_intrinsic_instr *instr)
case nir_intrinsic_get_buffer_size:
visit_get_buffer_size(ctx, instr);
break;
case nir_intrinsic_barrier: {
case nir_intrinsic_control_barrier: {
unsigned* bsize = ctx->program->info->cs.block_size;
unsigned workgroup_size = bsize[0] * bsize[1] * bsize[2];
if (workgroup_size > ctx->program->wave_size)

4
src/amd/llvm/ac_nir_to_llvm.c
View File

@@ -3555,7 +3555,7 @@ static void visit_intrinsic(struct ac_nir_context *ctx,
break;
case nir_intrinsic_memory_barrier_tcs_patch:
break;
case nir_intrinsic_barrier:
case nir_intrinsic_control_barrier:
ac_emit_barrier(&ctx->ac, ctx->stage);
break;
case nir_intrinsic_shared_atomic_add:
@@ -4919,7 +4919,7 @@ scan_tess_ctrl(nir_cf_node *cf_node, unsigned *upper_block_tf_writemask,
continue;
nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr);
if (intrin->intrinsic == nir_intrinsic_barrier) {
if (intrin->intrinsic == nir_intrinsic_control_barrier) {
/* If we find a barrier in nested control flow put this in the
* too hard basket. In GLSL this is not possible but it is in

2
src/amd/vulkan/radv_meta_fast_clear.c
View File

@@ -87,7 +87,7 @@ build_dcc_decompress_compute_shader(struct radv_device *dev)
nir_intrinsic_instr *membar = nir_intrinsic_instr_create(b.shader, nir_intrinsic_memory_barrier);
nir_builder_instr_insert(&b, &membar->instr);
nir_intrinsic_instr *bar = nir_intrinsic_instr_create(b.shader, nir_intrinsic_barrier);
nir_intrinsic_instr *bar = nir_intrinsic_instr_create(b.shader, nir_intrinsic_control_barrier);
nir_builder_instr_insert(&b, &bar->instr);
nir_ssa_def *outval = &tex->dest.ssa;

2
src/broadcom/compiler/nir_to_vir.c
View File

@@ -2258,7 +2258,7 @@ ntq_emit_intrinsic(struct v3d_compile *c, nir_intrinsic_instr *instr)
*/
break;
case nir_intrinsic_barrier:
case nir_intrinsic_control_barrier:
/* Emit a TSY op to get all invocations in the workgroup
* (actually supergroup) to block until the last invocation
* reaches the TSY op.

2
src/compiler/glsl/glsl_to_nir.cpp
View File

@@ -2714,7 +2714,7 @@ nir_visitor::visit(ir_barrier *)
}
nir_intrinsic_instr *instr =
nir_intrinsic_instr_create(this->shader, nir_intrinsic_barrier);
nir_intrinsic_instr_create(this->shader, nir_intrinsic_control_barrier);
nir_builder_instr_insert(&b, &instr->instr);
}

7
src/compiler/nir/nir_intrinsics.py
View File

@@ -195,7 +195,6 @@ intrinsic("get_buffer_size", src_comp=[-1], dest_comp=1,
def barrier(name):
intrinsic(name)
barrier("barrier")
barrier("discard")
# Demote fragment shader invocation to a helper invocation. Any stores to
@@ -207,6 +206,12 @@ barrier("discard")
barrier("demote")
intrinsic("is_helper_invocation", dest_comp=1, flags=[CAN_ELIMINATE])
# A workgroup-level control barrier. Any thread which hits this barrier will
# pause until all threads within the current workgroup have also hit the
# barrier. For compute shaders, the workgroup is defined as the local group.
# For tessellation control shaders, the workgroup is defined as the current
# patch. This intrinsic does not imply any sort of memory barrier.
barrier("control_barrier")
# Memory barrier with semantics analogous to the memoryBarrier() GLSL
# intrinsic.

2
src/compiler/nir/nir_opt_combine_stores.c
View File

@@ -303,7 +303,7 @@ combine_stores_block(struct combine_stores_state *state, nir_block *block)
update_combined_store(state, intrin);
break;
case nir_intrinsic_barrier:
case nir_intrinsic_control_barrier:
case nir_intrinsic_group_memory_barrier:
case nir_intrinsic_memory_barrier:
combine_stores_with_modes(state, nir_var_shader_out |

4
src/compiler/nir/nir_opt_copy_prop_vars.c
View File

@@ -164,7 +164,7 @@ gather_vars_written(struct copy_prop_var_state *state,
nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr);
switch (intrin->intrinsic) {
case nir_intrinsic_barrier:
case nir_intrinsic_control_barrier:
case nir_intrinsic_memory_barrier:
written->modes |= nir_var_shader_out |
nir_var_mem_ssbo |
@@ -798,7 +798,7 @@ copy_prop_vars_block(struct copy_prop_var_state *state,
nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr);
switch (intrin->intrinsic) {
case nir_intrinsic_barrier:
case nir_intrinsic_control_barrier:
case nir_intrinsic_memory_barrier:
if (debug) dump_instr(instr);

2
src/compiler/nir/nir_opt_dead_write_vars.c
View File

@@ -131,7 +131,7 @@ remove_dead_write_vars_local(void *mem_ctx, nir_block *block)
nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr);
switch (intrin->intrinsic) {
case nir_intrinsic_barrier:
case nir_intrinsic_control_barrier:
case nir_intrinsic_memory_barrier: {
clear_unused_for_modes(&unused_writes, nir_var_shader_out |
nir_var_mem_ssbo |

2
src/compiler/nir/nir_schedule.c
View File

@@ -347,7 +347,7 @@ nir_schedule_intrinsic_deps(nir_deps_state *state,
add_write_dep(state, &state->store_shared, n);
break;
case nir_intrinsic_barrier:
case nir_intrinsic_control_barrier:
case nir_intrinsic_memory_barrier_shared:
add_write_dep(state, &state->store_shared, n);

7
src/compiler/nir/tests/load_store_vectorizer_tests.cpp
View File

@@ -773,12 +773,13 @@ TEST_F(nir_load_store_vectorize_test, ssbo_load_adjacent_memory_barrier)
ASSERT_EQ(count_intrinsics(nir_intrinsic_load_ssbo), 2);
}
/* nir_intrinsic_barrier only syncs invocations in a workgroup, it doesn't
* require that loads/stores complete. */
/* nir_intrinsic_control_barrier only syncs invocations in a workgroup, it
* doesn't require that loads/stores complete.
*/
TEST_F(nir_load_store_vectorize_test, ssbo_load_adjacent_barrier)
{
create_load(nir_var_mem_ssbo, 0, 0, 0x1);
nir_builder_instr_insert(b, &nir_intrinsic_instr_create(b->shader, nir_intrinsic_barrier)->instr);
nir_builder_instr_insert(b, &nir_intrinsic_instr_create(b->shader, nir_intrinsic_control_barrier)->instr);
create_load(nir_var_mem_ssbo, 0, 4, 0x2);
nir_validate_shader(b->shader, NULL);

2
src/compiler/spirv/spirv_to_nir.c
View File

@@ -3712,7 +3712,7 @@ vtn_handle_barrier(struct vtn_builder *b, SpvOp opcode,
vtn_emit_memory_barrier(b, memory_scope, memory_semantics);
if (execution_scope == SpvScopeWorkgroup)
vtn_emit_barrier(b, nir_intrinsic_barrier);
vtn_emit_barrier(b, nir_intrinsic_control_barrier);
break;
}

4
src/freedreno/ir3/ir3_compiler_nir.c
View File

@@ -1155,7 +1155,7 @@ emit_intrinsic_barrier(struct ir3_context *ctx, nir_intrinsic_instr *intr)
struct ir3_instruction *barrier;
switch (intr->intrinsic) {
case nir_intrinsic_barrier:
case nir_intrinsic_control_barrier:
barrier = ir3_BAR(b);
barrier->cat7.g = true;
barrier->cat7.l = true;
@@ -1641,7 +1641,7 @@ emit_intrinsic(struct ir3_context *ctx, nir_intrinsic_instr *intr)
ctx->so->no_earlyz = true;
dst[0] = ctx->funcs->emit_intrinsic_atomic_image(ctx, intr);
break;
case nir_intrinsic_barrier:
case nir_intrinsic_control_barrier:
case nir_intrinsic_memory_barrier:
case nir_intrinsic_group_memory_barrier:
case nir_intrinsic_memory_barrier_atomic_counter:

2
src/freedreno/ir3/ir3_nir_lower_tess.c
View File

@@ -348,7 +348,7 @@ lower_tess_ctrl_block(nir_block *block, nir_builder *b, struct state *state)
nir_instr_remove(&intr->instr);
break;
case nir_intrinsic_barrier:
case nir_intrinsic_control_barrier:
case nir_intrinsic_memory_barrier_tcs_patch:
/* Hull shaders dispatch 32 wide so an entire patch will always
* fit in a single warp and execute in lock-step. Consequently,

2
src/gallium/auxiliary/gallivm/lp_bld_nir.c
View File

@@ -1352,7 +1352,7 @@ static void visit_intrinsic(struct lp_build_nir_context *bld_base,
case nir_intrinsic_shared_atomic_comp_swap:
visit_shared_atomic(bld_base, instr, result);
break;
case nir_intrinsic_barrier:
case nir_intrinsic_control_barrier:
visit_barrier(bld_base);
break;
case nir_intrinsic_memory_barrier:

2
src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
View File

@@ -2644,7 +2644,7 @@ Converter::visit(nir_intrinsic_instr *insn)
break;
}
case nir_intrinsic_barrier: {
case nir_intrinsic_control_barrier: {
// TODO: add flag to shader_info
info->numBarriers = 1;
Instruction *bar = mkOp2(OP_BAR, TYPE_U32, NULL, mkImm(0), mkImm(0));

4
src/intel/compiler/brw_fs_nir.cpp
View File

@@ -2751,7 +2751,7 @@ fs_visitor::nir_emit_tcs_intrinsic(const fs_builder &bld,
brw_imm_d(tcs_key->input_vertices));
break;
case nir_intrinsic_barrier: {
case nir_intrinsic_control_barrier: {
if (tcs_prog_data->instances == 1)
break;
@@ -3766,7 +3766,7 @@ fs_visitor::nir_emit_cs_intrinsic(const fs_builder &bld,
dest = get_nir_dest(instr->dest);
switch (instr->intrinsic) {
case nir_intrinsic_barrier:
case nir_intrinsic_control_barrier:
emit_barrier();
cs_prog_data->uses_barrier = true;
break;

2
src/intel/compiler/brw_vec4_tcs.cpp
View File

@@ -308,7 +308,7 @@ vec4_tcs_visitor::nir_emit_intrinsic(nir_intrinsic_instr *instr)
break;
}
case nir_intrinsic_barrier: {
case nir_intrinsic_control_barrier: {
dst_reg header = dst_reg(this, glsl_type::uvec4_type);
emit(TCS_OPCODE_CREATE_BARRIER_HEADER, header);
emit(SHADER_OPCODE_BARRIER, dst_null_ud(), src_reg(header));