diff options
-rw-r--r-- | src/amd/common/ac_llvm_helper.cpp | 3 | ||||
-rw-r--r-- | src/amd/common/ac_llvm_util.c | 2 | ||||
-rw-r--r-- | src/amd/common/ac_llvm_util.h | 1 | ||||
-rw-r--r-- | src/amd/common/ac_nir_to_llvm.c | 6 | ||||
-rw-r--r-- | src/gallium/auxiliary/gallivm/lp_bld_intr.c | 2 | ||||
-rw-r--r-- | src/gallium/auxiliary/gallivm/lp_bld_intr.h | 1 | ||||
-rw-r--r-- | src/gallium/drivers/radeonsi/si_shader.c | 17 |
7 files changed, 8 insertions, 24 deletions
diff --git a/src/amd/common/ac_llvm_helper.cpp b/src/amd/common/ac_llvm_helper.cpp index e42d00280bc..793737c395c 100644 --- a/src/amd/common/ac_llvm_helper.cpp +++ b/src/amd/common/ac_llvm_helper.cpp @@ -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) diff --git a/src/amd/common/ac_llvm_util.c b/src/amd/common/ac_llvm_util.c index 429904c0403..5fd785ad24a 100644 --- a/src/amd/common/ac_llvm_util.c +++ b/src/amd/common/ac_llvm_util.c @@ -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"; diff --git a/src/amd/common/ac_llvm_util.h b/src/amd/common/ac_llvm_util.h index 84fcbf111cf..29dc0c1c7d1 100644 --- a/src/amd/common/ac_llvm_util.h +++ b/src/amd/common/ac_llvm_util.h @@ -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), diff --git a/src/amd/common/ac_nir_to_llvm.c b/src/amd/common/ac_nir_to_llvm.c index 35f3c587228..bd7d77553e8 100644 --- a/src/amd/common/ac_nir_to_llvm.c +++ b/src/amd/common/ac_nir_to_llvm.c @@ -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) { diff --git a/src/gallium/auxiliary/gallivm/lp_bld_intr.c b/src/gallium/auxiliary/gallivm/lp_bld_intr.c index b92455593f4..74ed16f33f0 100644 --- a/src/gallium/auxiliary/gallivm/lp_bld_intr.c +++ b/src/gallium/auxiliary/gallivm/lp_bld_intr.c @@ -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"; diff --git a/src/gallium/auxiliary/gallivm/lp_bld_intr.h b/src/gallium/auxiliary/gallivm/lp_bld_intr.h index 0a929c51970..bf8143df87d 100644 --- a/src/gallium/auxiliary/gallivm/lp_bld_intr.h +++ b/src/gallium/auxiliary/gallivm/lp_bld_intr.h @@ -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), diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c index 8d3e34f91c8..787af9bae9a 100644 --- a/src/gallium/drivers/radeonsi/si_shader.c +++ b/src/gallium/drivers/radeonsi/si_shader.c @@ -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); |