ac: don't use byval LLVM qualifier in shaders
shader-db doesn't show any regression and 32-bit pointers with byval are declared as VGPRs for some reason. Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
This commit is contained in:
parent
0e40c6a7b7
commit
0d62370bbb
|
@ -60,8 +60,7 @@ bool ac_is_sgpr_param(LLVMValueRef arg)
|
|||
llvm::Argument *A = llvm::unwrap<llvm::Argument>(arg);
|
||||
llvm::AttributeList AS = A->getParent()->getAttributes();
|
||||
unsigned ArgNo = A->getArgNo();
|
||||
return AS.hasAttribute(ArgNo + 1, llvm::Attribute::ByVal) ||
|
||||
AS.hasAttribute(ArgNo + 1, llvm::Attribute::InReg);
|
||||
return AS.hasAttribute(ArgNo + 1, llvm::Attribute::InReg);
|
||||
}
|
||||
|
||||
LLVMValueRef ac_llvm_get_called_value(LLVMValueRef call)
|
||||
|
|
|
@ -152,7 +152,6 @@ static LLVMAttribute ac_attr_to_llvm_attr(enum ac_func_attr attr)
|
|||
{
|
||||
switch (attr) {
|
||||
case AC_FUNC_ATTR_ALWAYSINLINE: return LLVMAlwaysInlineAttribute;
|
||||
case AC_FUNC_ATTR_BYVAL: return LLVMByValAttribute;
|
||||
case AC_FUNC_ATTR_INREG: return LLVMInRegAttribute;
|
||||
case AC_FUNC_ATTR_NOALIAS: return LLVMNoAliasAttribute;
|
||||
case AC_FUNC_ATTR_NOUNWIND: return LLVMNoUnwindAttribute;
|
||||
|
@ -170,7 +169,6 @@ static const char *attr_to_str(enum ac_func_attr attr)
|
|||
{
|
||||
switch (attr) {
|
||||
case AC_FUNC_ATTR_ALWAYSINLINE: return "alwaysinline";
|
||||
case AC_FUNC_ATTR_BYVAL: return "byval";
|
||||
case AC_FUNC_ATTR_INREG: return "inreg";
|
||||
case AC_FUNC_ATTR_NOALIAS: return "noalias";
|
||||
case AC_FUNC_ATTR_NOUNWIND: return "nounwind";
|
||||
|
|
|
@ -37,7 +37,6 @@ extern "C" {
|
|||
|
||||
enum ac_func_attr {
|
||||
AC_FUNC_ATTR_ALWAYSINLINE = (1 << 0),
|
||||
AC_FUNC_ATTR_BYVAL = (1 << 1),
|
||||
AC_FUNC_ATTR_INREG = (1 << 2),
|
||||
AC_FUNC_ATTR_NOALIAS = (1 << 3),
|
||||
AC_FUNC_ATTR_NOUNWIND = (1 << 4),
|
||||
|
|
|
@ -323,15 +323,13 @@ create_llvm_function(LLVMContextRef ctx, LLVMModuleRef module,
|
|||
|
||||
LLVMSetFunctionCallConv(main_function, RADEON_LLVM_AMDGPU_CS);
|
||||
for (unsigned i = 0; i < args->sgpr_count; ++i) {
|
||||
ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_INREG);
|
||||
|
||||
if (args->array_params_mask & (1 << i)) {
|
||||
LLVMValueRef P = LLVMGetParam(main_function, i);
|
||||
ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_BYVAL);
|
||||
ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_NOALIAS);
|
||||
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) {
|
||||
|
|
|
@ -126,7 +126,6 @@ static LLVMAttribute lp_attr_to_llvm_attr(enum lp_func_attr attr)
|
|||
{
|
||||
switch (attr) {
|
||||
case LP_FUNC_ATTR_ALWAYSINLINE: return LLVMAlwaysInlineAttribute;
|
||||
case LP_FUNC_ATTR_BYVAL: return LLVMByValAttribute;
|
||||
case LP_FUNC_ATTR_INREG: return LLVMInRegAttribute;
|
||||
case LP_FUNC_ATTR_NOALIAS: return LLVMNoAliasAttribute;
|
||||
case LP_FUNC_ATTR_NOUNWIND: return LLVMNoUnwindAttribute;
|
||||
|
@ -144,7 +143,6 @@ static const char *attr_to_str(enum lp_func_attr attr)
|
|||
{
|
||||
switch (attr) {
|
||||
case LP_FUNC_ATTR_ALWAYSINLINE: return "alwaysinline";
|
||||
case LP_FUNC_ATTR_BYVAL: return "byval";
|
||||
case LP_FUNC_ATTR_INREG: return "inreg";
|
||||
case LP_FUNC_ATTR_NOALIAS: return "noalias";
|
||||
case LP_FUNC_ATTR_NOUNWIND: return "nounwind";
|
||||
|
|
|
@ -48,7 +48,6 @@
|
|||
|
||||
enum lp_func_attr {
|
||||
LP_FUNC_ATTR_ALWAYSINLINE = (1 << 0),
|
||||
LP_FUNC_ATTR_BYVAL = (1 << 1),
|
||||
LP_FUNC_ATTR_INREG = (1 << 2),
|
||||
LP_FUNC_ATTR_NOALIAS = (1 << 3),
|
||||
LP_FUNC_ATTR_NOUNWIND = (1 << 4),
|
||||
|
|
|
@ -4452,18 +4452,18 @@ static void si_create_function(struct si_shader_context *ctx,
|
|||
LLVMValueRef P = LLVMGetParam(ctx->main_fn, i);
|
||||
|
||||
/* The combination of:
|
||||
* - ByVal
|
||||
* - noalias
|
||||
* - dereferenceable
|
||||
* - invariant.load
|
||||
* allows the optimization passes to move loads and reduces
|
||||
* SGPR spilling significantly.
|
||||
*/
|
||||
lp_add_function_attr(ctx->main_fn, i + 1, LP_FUNC_ATTR_INREG);
|
||||
|
||||
if (LLVMGetTypeKind(LLVMTypeOf(P)) == LLVMPointerTypeKind) {
|
||||
lp_add_function_attr(ctx->main_fn, i + 1, LP_FUNC_ATTR_BYVAL);
|
||||
lp_add_function_attr(ctx->main_fn, i + 1, LP_FUNC_ATTR_NOALIAS);
|
||||
ac_add_attr_dereferenceable(P, UINT64_MAX);
|
||||
} else
|
||||
lp_add_function_attr(ctx->main_fn, i + 1, LP_FUNC_ATTR_INREG);
|
||||
}
|
||||
}
|
||||
|
||||
for (i = 0; i < fninfo->num_params; ++i) {
|
||||
|
@ -6595,15 +6595,8 @@ static void si_build_wrapper_function(struct si_shader_context *ctx,
|
|||
param_size = ac_get_type_size(param_type) / 4;
|
||||
is_sgpr = ac_is_sgpr_param(param);
|
||||
|
||||
if (is_sgpr) {
|
||||
#if HAVE_LLVM < 0x0400
|
||||
LLVMRemoveAttribute(param, LLVMByValAttribute);
|
||||
#else
|
||||
unsigned kind_id = LLVMGetEnumAttributeKindForName("byval", 5);
|
||||
LLVMRemoveEnumAttributeAtIndex(parts[part], param_idx + 1, kind_id);
|
||||
#endif
|
||||
if (is_sgpr)
|
||||
lp_add_function_attr(parts[part], param_idx + 1, LP_FUNC_ATTR_INREG);
|
||||
}
|
||||
|
||||
assert(out_idx + param_size <= (is_sgpr ? num_out_sgpr : num_out));
|
||||
assert(is_sgpr || out_idx >= num_out_sgpr);
|
||||
|
|
Loading…
Reference in New Issue