X-Git-Url: https://git.libre-soc.org/?a=blobdiff_plain;f=src%2Fgallium%2Fdrivers%2Fradeonsi%2Fsi_shader_llvm.c;h=11f5a14d121a9bd672b52c5e3fbb854ba4d571c0;hb=cf65c6f0d281bc64c9cea07b7f394036b2eeeaad;hp=64ceaf7ed34e4b494175ae1d04c2cb9f2337c1b6;hpb=420fe1e7f9ef56177c8f45e98e057488a2b57646;p=mesa.git

diff --git a/src/gallium/drivers/radeonsi/si_shader_llvm.c b/src/gallium/drivers/radeonsi/si_shader_llvm.c
index 64ceaf7ed34..11f5a14d121 100644
--- a/src/gallium/drivers/radeonsi/si_shader_llvm.c
+++ b/src/gallium/drivers/radeonsi/si_shader_llvm.c
@@ -114,8 +114,7 @@ void si_shader_binary_clean(struct si_shader_binary *binary)
 void si_llvm_context_init(struct si_shader_context *ctx,
 			  struct si_screen *sscreen,
 			  struct ac_llvm_compiler *compiler,
-			  unsigned wave_size,
-			  unsigned ballot_mask_bits)
+			  unsigned wave_size)
 {
 	/* Initialize the gallivm object:
 	 * We are only using the module, context, and builder fields of this struct.
@@ -129,7 +128,7 @@ void si_llvm_context_init(struct si_shader_context *ctx,
 	ac_llvm_context_init(&ctx->ac, compiler, sscreen->info.chip_class,
 			     sscreen->info.family,
 			     AC_FLOAT_MODE_NO_SIGNED_ZEROS_FP_MATH,
-			     wave_size, ballot_mask_bits);
+			     wave_size, 64);
 
 	ctx->voidt = LLVMVoidTypeInContext(ctx->ac.context);
 	ctx->i1 = LLVMInt1TypeInContext(ctx->ac.context);
@@ -149,13 +148,13 @@ void si_llvm_context_init(struct si_shader_context *ctx,
 	ctx->i1true = LLVMConstInt(ctx->i1, 1, 0);
 }
 
-/* Set the context to a certain TGSI shader. Can be called repeatedly
+/* Set the context to a certain shader. Can be called repeatedly
  * to change the shader. */
 void si_llvm_context_set_ir(struct si_shader_context *ctx,
 			    struct si_shader *shader)
 {
 	struct si_shader_selector *sel = shader->selector;
-	const struct tgsi_shader_info *info = &sel->info;
+	const struct si_shader_info *info = &sel->info;
 
 	ctx->shader = shader;
 	ctx->type = sel->type;
@@ -167,9 +166,9 @@ void si_llvm_context_set_ir(struct si_shader_context *ctx,
 	ctx->num_images = util_last_bit(info->images_declared);
 }
 
-void si_llvm_create_func(struct si_shader_context *ctx,
-			 const char *name,
-			 LLVMTypeRef *return_types, unsigned num_return_elems)
+void si_llvm_create_func(struct si_shader_context *ctx, const char *name,
+			 LLVMTypeRef *return_types, unsigned num_return_elems,
+			 unsigned max_workgroup_size)
 {
 	LLVMTypeRef ret_type;
 	enum ac_llvm_calling_convention call_conv;
@@ -217,6 +216,19 @@ void si_llvm_create_func(struct si_shader_context *ctx,
 	ctx->return_type = ret_type;
 	ctx->main_fn = ac_build_main(&ctx->args, &ctx->ac, call_conv, name,
 				     ret_type, ctx->ac.module);
+	ctx->return_value = LLVMGetUndef(ctx->return_type);
+
+	if (ctx->screen->info.address32_hi) {
+		ac_llvm_add_target_dep_function_attr(ctx->main_fn,
+						     "amdgpu-32bit-address-high-bits",
+						     ctx->screen->info.address32_hi);
+	}
+
+	LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
+					   "no-signed-zeros-fp-math",
+					   "true");
+
+	ac_llvm_set_workgroup_size(ctx->main_fn, max_workgroup_size);
 }
 
 void si_llvm_optimize_module(struct si_shader_context *ctx)