4 files changed, 3 insertions, 9 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) {