[Mesa-dev] [PATCH 2/2] ac/nir: set workgroup size attribute to correct value.

[Dave Airlie]

From: Dave Airlie <airlied at redhat.com>

This ports: 55445ff1891724c78e6573d2f8c721e14c0449fc from radeonsi

    radeonsi: tell LLVM not to remove s_barrier instructions

    LLVM 5.0 removes s_barrier instructions if the max-work-group-size
    attribute is not set. What a surprise.

Signed-off-by: Dave Airlie <airlied at redhat.com>
---
 src/amd/common/ac_nir_to_llvm.c | 43 ++++++++++++++++++++++++++++++++++++++---
 1 file changed, 40 insertions(+), 3 deletions(-)

diff --git a/src/amd/common/ac_nir_to_llvm.c b/src/amd/common/ac_nir_to_llvm.c
index 28ba47d..7f37178 100644
--- a/src/amd/common/ac_nir_to_llvm.c
+++ b/src/amd/common/ac_nir_to_llvm.c
@@ -33,6 +33,8 @@
 #include "ac_shader_info.h"
 #include "ac_exp_param.h"
 
+#define AC_SI_MAX_VARIABLE_THREADS_PER_BLOCK 1024
+
 enum radeon_llvm_calling_convention {
 	RADEON_LLVM_AMDGPU_VS = 87,
 	RADEON_LLVM_AMDGPU_GS = 88,
@@ -57,7 +59,7 @@ struct nir_to_llvm_context {
 	struct ac_llvm_context ac;
 	const struct ac_nir_compiler_options *options;
 	struct ac_shader_variant_info *shader_info;
-
+	unsigned max_workgroup_size;
 	LLVMContextRef context;
 	LLVMModuleRef module;
 	LLVMBuilderRef builder;
@@ -257,7 +259,8 @@ create_llvm_function(LLVMContextRef ctx, LLVMModuleRef module,
                      LLVMBuilderRef builder, LLVMTypeRef *return_types,
                      unsigned num_return_elems, LLVMTypeRef *param_types,
                      unsigned param_count, unsigned array_params_mask,
-                     unsigned sgpr_params, bool unsafe_math)
+                     unsigned sgpr_params, unsigned max_workgroup_size,
+		     bool unsafe_math)
 {
 	LLVMTypeRef main_function_type, ret_type;
 	LLVMBasicBlockRef main_function_body;
@@ -289,6 +292,11 @@ create_llvm_function(LLVMContextRef ctx, LLVMModuleRef module,
 		}
 	}
 
+	if (max_workgroup_size) {
+		ac_llvm_add_target_dep_function_attr(main_function,
+						     "amdgpu-max-work-group-size",
+						     max_workgroup_size);
+	}
 	if (unsafe_math) {
 		/* These were copied from some LLVM test. */
 		LLVMAddTargetDependentFunctionAttr(main_function,
@@ -773,7 +781,8 @@ static void create_function(struct nir_to_llvm_context *ctx)
 
 	ctx->main_function = create_llvm_function(
 	    ctx->context, ctx->module, ctx->builder, NULL, 0, arg_types,
-	    arg_idx, array_params_mask, sgpr_count, ctx->options->unsafe_math);
+	    arg_idx, array_params_mask, sgpr_count, ctx->max_workgroup_size,
+	    ctx->options->unsafe_math);
 	set_llvm_calling_convention(ctx->main_function, ctx->stage);
 
 	ctx->shader_info->num_input_sgprs = 0;
@@ -5855,6 +5864,33 @@ ac_setup_rings(struct nir_to_llvm_context *ctx)
 	}
 }
 
+static unsigned
+ac_nir_get_max_workgroup_size(enum chip_class chip_class,
+			      struct nir_shader *nir)
+{
+	switch (nir->stage) {
+	case MESA_SHADER_TESS_CTRL:
+		return chip_class >= CIK ? 128 : 64;
+	case MESA_SHADER_GEOMETRY:
+		return 64;
+	case MESA_SHADER_COMPUTE:
+		break;
+	default:
+		return 0;
+	}
+
+	unsigned max_workgroup_size = nir->info.cs.local_size[0] *
+		nir->info.cs.local_size[1] *
+		nir->info.cs.local_size[2];
+	if (!max_workgroup_size) {
+		/* This is a variable group size compute shader,
+		 * compile it for the maximum possible group size.
+		 */
+		max_workgroup_size = AC_SI_MAX_VARIABLE_THREADS_PER_BLOCK;
+	}
+	return max_workgroup_size;
+}
+
 static
 LLVMModuleRef ac_translate_nir_to_llvm(LLVMTargetMachineRef tm,
                                        struct nir_shader *nir,
@@ -5891,6 +5927,7 @@ LLVMModuleRef ac_translate_nir_to_llvm(LLVMTargetMachineRef tm,
 	ctx.builder = LLVMCreateBuilderInContext(ctx.context);
 	ctx.ac.builder = ctx.builder;
 	ctx.stage = nir->stage;
+	ctx.max_workgroup_size = ac_nir_get_max_workgroup_size(ctx.options->chip_class, nir);
 
 	for (i = 0; i < AC_UD_MAX_SETS; i++)
 		shader_info->user_sgprs_locs.descriptor_sets[i].sgpr_idx = -1;
-- 
2.9.3