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>

diff --git a/src/amd/common/ac_llvm_helper.cpp b/src/amd/common/ac_llvm_helper.cpp
index e42d00280bc71bc24f05d99bd0ada169308b9dbf..793737c395c528f9d2c9f3af6f01ab4f18aeac63 100644 (file)
--- 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 429904c04036399366f2029398e269a3aa24d883..5fd785ad24ad40317e98c3c6e30b49fa6ef09e84 100644 (file)
--- 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 84fcbf111cf125fc0aca01e7daf1db077ed5b2d7..29dc0c1c7d1d1de0ebb7c0ff95a587ed2ea87ab2 100644 (file)
--- 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 35f3c587228bac33705c59baedc8fb68b7a3a0ca..bd7d77553e8fa8e90dd6f2cd96abcf4ffd8b6bba 100644 (file)
--- 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 b92455593f4d611758b4322dcb55df468c5fa2a5..74ed16f33f0e5dae13a9587cbe36e236a9a0c6b8 100644 (file)
--- 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 0a929c51970e8721e7329e6ff207ea9876d1be46..bf8143df87db4c35a76199fc70b5f5d12ce40057 100644 (file)
--- 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 8d3e34f91c894ecbd4aa92bd075d263f53cd7598..787af9bae9a35820263e371a2143b11a8119b883 100644 (file)
--- 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);