Mesa (master): ac: don't use byval LLVM qualifier in shaders
Marek Olšák
mareko at kemper.freedesktop.org
Sat Jan 27 01:09:28 UTC 2018
Module: Mesa
Branch: master
Commit: 0d62370bbb9a70bc4d493fa8be9ddf73c87d15d9
URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=0d62370bbb9a70bc4d493fa8be9ddf73c87d15d9
Author: Marek Olšák <marek.olsak at amd.com>
Date: Mon Jan 1 00:30:51 2018 +0100
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 at gmail.com>
---
src/amd/common/ac_llvm_helper.cpp | 3 +--
src/amd/common/ac_llvm_util.c | 2 --
src/amd/common/ac_llvm_util.h | 1 -
src/amd/common/ac_nir_to_llvm.c | 6 ++----
src/gallium/auxiliary/gallivm/lp_bld_intr.c | 2 --
src/gallium/auxiliary/gallivm/lp_bld_intr.h | 1 -
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 e42d00280b..793737c395 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 429904c040..5fd785ad24 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 84fcbf111c..29dc0c1c7d 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 35f3c58722..bd7d77553e 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 b92455593f..74ed16f33f 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 0a929c5197..bf8143df87 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 8d3e34f91c..787af9bae9 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);
More information about the mesa-commit
mailing list