[Mesa-dev] [PATCH 14/16] ac: use amdgpu-flat-work-group-size

[Marek Olšák]

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

the old one is being deprecated or removed
---
 src/amd/common/ac_nir_to_llvm.c          | 3 ++-
 src/gallium/drivers/radeonsi/si_shader.c | 4 +++-
 2 files changed, 5 insertions(+), 2 deletions(-)

diff --git a/src/amd/common/ac_nir_to_llvm.c b/src/amd/common/ac_nir_to_llvm.c
index 11ba487..4492d8e 100644
--- a/src/amd/common/ac_nir_to_llvm.c
+++ b/src/amd/common/ac_nir_to_llvm.c
@@ -346,21 +346,22 @@ create_llvm_function(LLVMContextRef ctx, LLVMModuleRef module,
 			ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_BYVAL);
 			ac_add_attr_dereferenceable(P, UINT64_MAX);
 		}
 		else {
 			ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_INREG);
 		}
 	}
 
 	if (max_workgroup_size) {
 		ac_llvm_add_target_dep_function_attr(main_function,
-						     "amdgpu-max-work-group-size",
+						     HAVE_LLVM >= 0x0400 ? "amdgpu-max-work-group-size" :
+								           "amdgpu-flat-work-group-size",
 						     max_workgroup_size);
 	}
 	if (unsafe_math) {
 		/* These were copied from some LLVM test. */
 		LLVMAddTargetDependentFunctionAttr(main_function,
 						   "less-precise-fpmad",
 						   "true");
 		LLVMAddTargetDependentFunctionAttr(main_function,
 						   "no-infs-fp-math",
 						   "true");
diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c
index ff372ae..506da6f 100644
--- a/src/gallium/drivers/radeonsi/si_shader.c
+++ b/src/gallium/drivers/radeonsi/si_shader.c
@@ -4155,21 +4155,23 @@ static void si_create_function(struct si_shader_context *ctx,
 		} else
 			lp_add_function_attr(ctx->main_fn, i + 1, LP_FUNC_ATTR_INREG);
 	}
 
 	for (i = 0; i < fninfo->num_params; ++i) {
 		if (fninfo->assign[i])
 			*fninfo->assign[i] = LLVMGetParam(ctx->main_fn, i);
 	}
 
 	if (max_workgroup_size) {
-		si_llvm_add_attribute(ctx->main_fn, "amdgpu-max-work-group-size",
+		si_llvm_add_attribute(ctx->main_fn,
+				      HAVE_LLVM >= 0x0400 ? "amdgpu-max-work-group-size" :
+							    "amdgpu-flat-work-group-size",
 				      max_workgroup_size);
 	}
 	LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
 					   "no-signed-zeros-fp-math",
 					   "true");
 
 	if (ctx->screen->b.debug_flags & DBG(UNSAFE_MATH)) {
 		/* These were copied from some LLVM test. */
 		LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
 						   "less-precise-fpmad",
-- 
2.7.4