ac: use amdgpu-flat-work-group-size

Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>

src/amd/common/ac_llvm_util.c

@@ -269,6 +269,16 @@ ac_llvm_add_target_dep_function_attr(LLVMValueRef F,
 	LLVMAddTargetDependentFunctionAttr(F, name, str);
 }
 
+void ac_llvm_set_workgroup_size(LLVMValueRef F, unsigned size)
+{
+	if (!size)
+		return;
+
+	char str[32];
+	snprintf(str, sizeof(str), "%u,%u", size, size);
+	LLVMAddTargetDependentFunctionAttr(F, "amdgpu-flat-work-group-size", str);
+}
+
 unsigned
 ac_count_scratch_private_memory(LLVMValueRef function)
 {

src/amd/common/ac_llvm_util.h

@@ -109,6 +109,7 @@ LLVMBuilderRef ac_create_builder(LLVMContextRef ctx,
 void
 ac_llvm_add_target_dep_function_attr(LLVMValueRef F,
 				     const char *name, unsigned value);
+void ac_llvm_set_workgroup_size(LLVMValueRef F, unsigned size);
 
 static inline unsigned
 ac_get_load_intr_attribs(bool can_speculate)

src/amd/vulkan/radv_nir_to_llvm.c

@@ -518,11 +518,8 @@ create_llvm_function(LLVMContextRef ctx, LLVMModuleRef module,
 						     options->address32_hi);
 	}
 
-	if (max_workgroup_size) {
-		ac_llvm_add_target_dep_function_attr(main_function,
-						     "amdgpu-max-work-group-size",
-						     max_workgroup_size);
-	}
+	ac_llvm_set_workgroup_size(main_function, max_workgroup_size);
+
 	if (options->unsafe_math) {
 		/* These were copied from some LLVM test. */
 		LLVMAddTargetDependentFunctionAttr(main_function,

src/gallium/drivers/radeonsi/si_shader.c

@@ -4307,11 +4307,8 @@ void si_create_function(struct si_shader_context *ctx,
 						     ctx->screen->info.address32_hi);
 	}
 
-	if (max_workgroup_size) {
-		ac_llvm_add_target_dep_function_attr(ctx->main_fn,
-						     "amdgpu-max-work-group-size",
-						     max_workgroup_size);
-	}
+	ac_llvm_set_workgroup_size(ctx->main_fn, max_workgroup_size);
+
 	LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
 					   "no-signed-zeros-fp-math",
 					   "true");