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 <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3307>

diff --git a/src/amd/compiler/aco_instruction_selection.cpp b/src/amd/compiler/aco_instruction_selection.cpp
index 29dea6e6cd35f1f7b319cc0d99df5067de9be331..a106631b4c1e4955ee74da1642d98e5a0733742c 100644 (file)
--- 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 fb042ac1d4adcb59b7cf0b7f3d7b2fbe3f774e95..17c21cadad9129782dc57aeb136cd25ef4b4ae31 100644 (file)
--- 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 e0e83c2754ff7595e0c3c7147d26dbd91bd242f6..8e54c7286aa31192f42159602268afe1cb7abb72 100644 (file)
--- 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 401958e1471dcba7101d3830dab2d5214f181c87..ffd166f90f462690d534eed1fab55c7d51ee6107 100644 (file)
--- 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 4ab620142b07ca9efb710aeaa7ba31382b296c5f..a3a85786fcd35095c4dbc10b9f77bf598f24da1b 100644 (file)
--- 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 04f58b0172bdc6856107605e6fe35f0aef11c32f..3768016a78787dd508b2fc3e8a2f76033b859ee9 100644 (file)
--- 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 508833b5b4683fcebce704483e826c3dc158bf5c..791ec094c00563737e084c809b2c530d569b5c05 100644 (file)
--- 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 26ca61969aff50f7ed5d632e9f57be2f6c89253a..dbce55bfb8ca1fa852827d1833123321e7340a52 100644 (file)
--- 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 ecbe77895d3124bf9b85ad5e4a82e9b8fb4c2e8f..578e8706818e9dd79d38ccbbc49f61c7db3b93bc 100644 (file)
--- 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 0ad95a5bc89588f05d1ce0a5035239bb152d9724..8428f867cccd306c7418303f97cd3628b317667c 100644 (file)
--- 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 0b332d8bc78bacd5cc24c8d24b11cc15d24c197e..cd623704fc38323c30943fada9daa4caace2a659 100644 (file)
--- 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 2ed2bf34875f65538d0da8b5d05c2a8067383d8b..0a0df6ba8470d6998f532ae882acb830676ab05e 100644 (file)
--- 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 133388a3a80a9df221467e4f322be9268c779eb3..f1ba10456ed78ad6592c982c9f898bdca6cba3ea 100644 (file)
--- 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 b29903537b8823177fb2a8f5e3c984edac46150d..056b009ef752f38ec1dfdea06cb3898a58770b77 100644 (file)
--- 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 72a2721e0304f22f488840fb970aea9cf3d26a2b..34e6420c58802af9427a1f4ef822640ab9db2391 100644 (file)
--- 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 13b60dd4e8636ad16e0b314fb05fe61fc177bbcf..8a0274e30a6c3557d04c1a67f19c2e8dd3420d79 100644 (file)
--- 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 a104df497cd4b0e90706d99ec3f52d24fd50d879..c794bb90d978be682e5263565acecb3e363bd581 100644 (file)
--- 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 e539657335ca1d91ac2312e4c3bcdc8188c178f7..d9df1fcfaf76419468bf21133fbee8d0113e3f1e 100644 (file)
--- 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));