From e40b11bbcb02dde1a8f989ca6545e22414c6f4ce Mon Sep 17 00:00:00 2001 From: Jason Ekstrand Date: Tue, 7 Jan 2020 14:54:26 -0600 Subject: [PATCH] nir: Rename nir_intrinsic_barrier to control_barrier 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 Part-of: --- src/amd/compiler/aco_instruction_selection.cpp | 2 +- src/amd/llvm/ac_nir_to_llvm.c | 4 ++-- src/amd/vulkan/radv_meta_fast_clear.c | 2 +- src/broadcom/compiler/nir_to_vir.c | 2 +- src/compiler/glsl/glsl_to_nir.cpp | 2 +- src/compiler/nir/nir_intrinsics.py | 7 ++++++- src/compiler/nir/nir_opt_combine_stores.c | 2 +- src/compiler/nir/nir_opt_copy_prop_vars.c | 4 ++-- src/compiler/nir/nir_opt_dead_write_vars.c | 2 +- src/compiler/nir/nir_schedule.c | 2 +- src/compiler/nir/tests/load_store_vectorizer_tests.cpp | 7 ++++--- src/compiler/spirv/spirv_to_nir.c | 2 +- src/freedreno/ir3/ir3_compiler_nir.c | 4 ++-- src/freedreno/ir3/ir3_nir_lower_tess.c | 2 +- src/gallium/auxiliary/gallivm/lp_bld_nir.c | 2 +- src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp | 2 +- src/intel/compiler/brw_fs_nir.cpp | 4 ++-- src/intel/compiler/brw_vec4_tcs.cpp | 2 +- 18 files changed, 30 insertions(+), 24 deletions(-) diff --git a/src/amd/compiler/aco_instruction_selection.cpp b/src/amd/compiler/aco_instruction_selection.cpp index 29dea6e6cd3..a106631b4c1 100644 --- a/src/amd/compiler/aco_instruction_selection.cpp +++ b/src/amd/compiler/aco_instruction_selection.cpp @@ -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) diff --git a/src/amd/llvm/ac_nir_to_llvm.c b/src/amd/llvm/ac_nir_to_llvm.c index fb042ac1d4a..17c21cadad9 100644 --- a/src/amd/llvm/ac_nir_to_llvm.c +++ b/src/amd/llvm/ac_nir_to_llvm.c @@ -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 diff --git a/src/amd/vulkan/radv_meta_fast_clear.c b/src/amd/vulkan/radv_meta_fast_clear.c index e0e83c2754f..8e54c7286aa 100644 --- a/src/amd/vulkan/radv_meta_fast_clear.c +++ b/src/amd/vulkan/radv_meta_fast_clear.c @@ -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; diff --git a/src/broadcom/compiler/nir_to_vir.c b/src/broadcom/compiler/nir_to_vir.c index 401958e1471..ffd166f90f4 100644 --- a/src/broadcom/compiler/nir_to_vir.c +++ b/src/broadcom/compiler/nir_to_vir.c @@ -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. diff --git a/src/compiler/glsl/glsl_to_nir.cpp b/src/compiler/glsl/glsl_to_nir.cpp index 4ab620142b0..a3a85786fcd 100644 --- a/src/compiler/glsl/glsl_to_nir.cpp +++ b/src/compiler/glsl/glsl_to_nir.cpp @@ -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); } diff --git a/src/compiler/nir/nir_intrinsics.py b/src/compiler/nir/nir_intrinsics.py index 04f58b0172b..3768016a787 100644 --- a/src/compiler/nir/nir_intrinsics.py +++ b/src/compiler/nir/nir_intrinsics.py @@ -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. diff --git a/src/compiler/nir/nir_opt_combine_stores.c b/src/compiler/nir/nir_opt_combine_stores.c index 508833b5b46..791ec094c00 100644 --- a/src/compiler/nir/nir_opt_combine_stores.c +++ b/src/compiler/nir/nir_opt_combine_stores.c @@ -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 | diff --git a/src/compiler/nir/nir_opt_copy_prop_vars.c b/src/compiler/nir/nir_opt_copy_prop_vars.c index 26ca61969af..dbce55bfb8c 100644 --- a/src/compiler/nir/nir_opt_copy_prop_vars.c +++ b/src/compiler/nir/nir_opt_copy_prop_vars.c @@ -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); diff --git a/src/compiler/nir/nir_opt_dead_write_vars.c b/src/compiler/nir/nir_opt_dead_write_vars.c index ecbe77895d3..578e8706818 100644 --- a/src/compiler/nir/nir_opt_dead_write_vars.c +++ b/src/compiler/nir/nir_opt_dead_write_vars.c @@ -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 | diff --git a/src/compiler/nir/nir_schedule.c b/src/compiler/nir/nir_schedule.c index 0ad95a5bc89..8428f867ccc 100644 --- a/src/compiler/nir/nir_schedule.c +++ b/src/compiler/nir/nir_schedule.c @@ -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); diff --git a/src/compiler/nir/tests/load_store_vectorizer_tests.cpp b/src/compiler/nir/tests/load_store_vectorizer_tests.cpp index 0b332d8bc78..cd623704fc3 100644 --- a/src/compiler/nir/tests/load_store_vectorizer_tests.cpp +++ b/src/compiler/nir/tests/load_store_vectorizer_tests.cpp @@ -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); diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c index 2ed2bf34875..0a0df6ba847 100644 --- a/src/compiler/spirv/spirv_to_nir.c +++ b/src/compiler/spirv/spirv_to_nir.c @@ -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; } diff --git a/src/freedreno/ir3/ir3_compiler_nir.c b/src/freedreno/ir3/ir3_compiler_nir.c index 133388a3a80..f1ba10456ed 100644 --- a/src/freedreno/ir3/ir3_compiler_nir.c +++ b/src/freedreno/ir3/ir3_compiler_nir.c @@ -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: diff --git a/src/freedreno/ir3/ir3_nir_lower_tess.c b/src/freedreno/ir3/ir3_nir_lower_tess.c index b29903537b8..056b009ef75 100644 --- a/src/freedreno/ir3/ir3_nir_lower_tess.c +++ b/src/freedreno/ir3/ir3_nir_lower_tess.c @@ -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, diff --git a/src/gallium/auxiliary/gallivm/lp_bld_nir.c b/src/gallium/auxiliary/gallivm/lp_bld_nir.c index 72a2721e030..34e6420c588 100644 --- a/src/gallium/auxiliary/gallivm/lp_bld_nir.c +++ b/src/gallium/auxiliary/gallivm/lp_bld_nir.c @@ -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: diff --git a/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp b/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp index 13b60dd4e86..8a0274e30a6 100644 --- a/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp +++ b/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp @@ -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)); diff --git a/src/intel/compiler/brw_fs_nir.cpp b/src/intel/compiler/brw_fs_nir.cpp index a104df497cd..c794bb90d97 100644 --- a/src/intel/compiler/brw_fs_nir.cpp +++ b/src/intel/compiler/brw_fs_nir.cpp @@ -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; diff --git a/src/intel/compiler/brw_vec4_tcs.cpp b/src/intel/compiler/brw_vec4_tcs.cpp index e539657335c..d9df1fcfaf7 100644 --- a/src/intel/compiler/brw_vec4_tcs.cpp +++ b/src/intel/compiler/brw_vec4_tcs.cpp @@ -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)); -- 2.30.2