[Mesa-dev] [PATCH 5/9] ac: add ac_build_s_barrier

[Marek Olšák]

From: Marek Olšák <marek.olsak at amd.com>

---
 src/amd/common/ac_llvm_build.c           | 6 ++++++
 src/amd/common/ac_llvm_build.h           | 1 +
 src/amd/common/ac_nir_to_llvm.c          | 3 +--
 src/gallium/drivers/radeonsi/si_shader.c | 4 +---
 4 files changed, 9 insertions(+), 5 deletions(-)

diff --git a/src/amd/common/ac_llvm_build.c b/src/amd/common/ac_llvm_build.c
index 87e36df6431..c89bdf49faf 100644
--- a/src/amd/common/ac_llvm_build.c
+++ b/src/amd/common/ac_llvm_build.c
@@ -341,20 +341,26 @@ void ac_build_type_name_for_intr(LLVMTypeRef type, char *buf, unsigned bufsize)
 LLVMValueRef
 ac_build_phi(struct ac_llvm_context *ctx, LLVMTypeRef type,
 	     unsigned count_incoming, LLVMValueRef *values,
 	     LLVMBasicBlockRef *blocks)
 {
 	LLVMValueRef phi = LLVMBuildPhi(ctx->builder, type, "");
 	LLVMAddIncoming(phi, values, blocks, count_incoming);
 	return phi;
 }
 
+void ac_build_s_barrier(struct ac_llvm_context *ctx)
+{
+	ac_build_intrinsic(ctx, "llvm.amdgcn.s.barrier", ctx->voidt, NULL,
+			   0, AC_FUNC_ATTR_CONVERGENT);
+}
+
 /* Prevent optimizations (at least of memory accesses) across the current
  * point in the program by emitting empty inline assembly that is marked as
  * having side effects.
  *
  * Optionally, a value can be passed through the inline assembly to prevent
  * LLVM from hoisting calls to ReadNone functions.
  */
 void
 ac_build_optimization_barrier(struct ac_llvm_context *ctx,
 			      LLVMValueRef *pvgpr)
diff --git a/src/amd/common/ac_llvm_build.h b/src/amd/common/ac_llvm_build.h
index c5753037e7b..84212f0d459 100644
--- a/src/amd/common/ac_llvm_build.h
+++ b/src/amd/common/ac_llvm_build.h
@@ -126,20 +126,21 @@ ac_build_intrinsic(struct ac_llvm_context *ctx, const char *name,
 		   LLVMTypeRef return_type, LLVMValueRef *params,
 		   unsigned param_count, unsigned attrib_mask);
 
 void ac_build_type_name_for_intr(LLVMTypeRef type, char *buf, unsigned bufsize);
 
 LLVMValueRef
 ac_build_phi(struct ac_llvm_context *ctx, LLVMTypeRef type,
 	     unsigned count_incoming, LLVMValueRef *values,
 	     LLVMBasicBlockRef *blocks);
 
+void ac_build_s_barrier(struct ac_llvm_context *ctx);
 void ac_build_optimization_barrier(struct ac_llvm_context *ctx,
 				   LLVMValueRef *pvgpr);
 
 LLVMValueRef ac_build_shader_clock(struct ac_llvm_context *ctx);
 
 LLVMValueRef ac_build_ballot(struct ac_llvm_context *ctx, LLVMValueRef value);
 
 LLVMValueRef ac_build_vote_all(struct ac_llvm_context *ctx, LLVMValueRef value);
 
 LLVMValueRef ac_build_vote_any(struct ac_llvm_context *ctx, LLVMValueRef value);
diff --git a/src/amd/common/ac_nir_to_llvm.c b/src/amd/common/ac_nir_to_llvm.c
index cffc980e51f..1584fef7ab7 100644
--- a/src/amd/common/ac_nir_to_llvm.c
+++ b/src/amd/common/ac_nir_to_llvm.c
@@ -2575,22 +2575,21 @@ static void emit_membar(struct ac_llvm_context *ac,
 void ac_emit_barrier(struct ac_llvm_context *ac, gl_shader_stage stage)
 {
 	/* SI only (thanks to a hw bug workaround):
 	 * The real barrier instruction isn’t needed, because an entire patch
 	 * always fits into a single wave.
 	 */
 	if (ac->chip_class == SI && stage == MESA_SHADER_TESS_CTRL) {
 		ac_build_waitcnt(ac, LGKM_CNT & VM_CNT);
 		return;
 	}
-	ac_build_intrinsic(ac, "llvm.amdgcn.s.barrier",
-			   ac->voidt, NULL, 0, AC_FUNC_ATTR_CONVERGENT);
+	ac_build_s_barrier(ac);
 }
 
 static void emit_discard(struct ac_nir_context *ctx,
 			 const nir_intrinsic_instr *instr)
 {
 	LLVMValueRef cond;
 
 	if (instr->intrinsic == nir_intrinsic_discard_if) {
 		cond = LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ,
 				     get_src(ctx, instr->src[0]),
diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c
index 81c825db1e4..29523474735 100644
--- a/src/gallium/drivers/radeonsi/si_shader.c
+++ b/src/gallium/drivers/radeonsi/si_shader.c
@@ -4397,23 +4397,21 @@ static void si_llvm_emit_barrier(const struct lp_build_tgsi_action *action,
 	/* SI only (thanks to a hw bug workaround):
 	 * The real barrier instruction isn’t needed, because an entire patch
 	 * always fits into a single wave.
 	 */
 	if (ctx->screen->info.chip_class == SI &&
 	    ctx->type == PIPE_SHADER_TESS_CTRL) {
 		ac_build_waitcnt(&ctx->ac, LGKM_CNT & VM_CNT);
 		return;
 	}
 
-	ac_build_intrinsic(&ctx->ac,
-			   "llvm.amdgcn.s.barrier",
-			   ctx->voidt, NULL, 0, AC_FUNC_ATTR_CONVERGENT);
+	ac_build_s_barrier(&ctx->ac);
 }
 
 static void si_create_function(struct si_shader_context *ctx,
 			       const char *name,
 			       LLVMTypeRef *returns, unsigned num_returns,
 			       struct si_function_info *fninfo,
 			       unsigned max_workgroup_size)
 {
 	int i;
 
-- 
2.17.1