[Mesa-dev] [PATCH 7/8] radeonsi: fix passing address32_hi to LLVM for high values

[Marek Olšák]

From: Marek Olšák <marek.ol...@amd.com>

The old function treats high values as negative, which LLVM interprets as 0.
---
 src/amd/common/ac_llvm_util.c            | 4 ++--
 src/amd/common/ac_llvm_util.h            | 2 +-
 src/gallium/drivers/radeonsi/si_shader.c | 7 +++++--
 3 files changed, 8 insertions(+), 5 deletions(-)

diff --git a/src/amd/common/ac_llvm_util.c b/src/amd/common/ac_llvm_util.c
index b88c4e4..16b5ec3 100644
--- a/src/amd/common/ac_llvm_util.c
+++ b/src/amd/common/ac_llvm_util.c
@@ -193,17 +193,17 @@ void ac_add_func_attributes(LLVMContextRef ctx, 
LLVMValueRef function,
 void
 ac_dump_module(LLVMModuleRef module)
 {
        char *str = LLVMPrintModuleToString(module);
        fprintf(stderr, "%s", str);
        LLVMDisposeMessage(str);
 }
 
 void
 ac_llvm_add_target_dep_function_attr(LLVMValueRef F,
-                                    const char *name, int value)
+                                    const char *name, unsigned value)
 {
        char str[16];
 
-       snprintf(str, sizeof(str), "%i", value);
+       snprintf(str, sizeof(str), "0x%x", value);
        LLVMAddTargetDependentFunctionAttr(F, name, str);
 }
diff --git a/src/amd/common/ac_llvm_util.h b/src/amd/common/ac_llvm_util.h
index 3cf385a..3229524 100644
--- a/src/amd/common/ac_llvm_util.h
+++ b/src/amd/common/ac_llvm_util.h
@@ -80,21 +80,21 @@ void ac_add_func_attributes(LLVMContextRef ctx, 
LLVMValueRef function,
 void ac_dump_module(LLVMModuleRef module);
 
 LLVMValueRef ac_llvm_get_called_value(LLVMValueRef call);
 bool ac_llvm_is_function(LLVMValueRef v);
 
 LLVMBuilderRef ac_create_builder(LLVMContextRef ctx,
                                 enum ac_float_mode float_mode);
 
 void
 ac_llvm_add_target_dep_function_attr(LLVMValueRef F,
-                                    const char *name, int value);
+                                    const char *name, unsigned value);
 
 static inline unsigned
 ac_get_load_intr_attribs(bool can_speculate)
 {
        /* READNONE means writes can't affect it, while READONLY means that
         * writes can affect it. */
        return can_speculate ? AC_FUNC_ATTR_READNONE :
                               AC_FUNC_ATTR_READONLY;
 }
 
diff --git a/src/gallium/drivers/radeonsi/si_shader.c 
b/src/gallium/drivers/radeonsi/si_shader.c
index 6be1806..60fc564 100644
--- a/src/gallium/drivers/radeonsi/si_shader.c
+++ b/src/gallium/drivers/radeonsi/si_shader.c
@@ -4411,22 +4411,25 @@ static void si_create_function(struct si_shader_context 
*ctx,
                        lp_add_function_attr(ctx->main_fn, i + 1, 
LP_FUNC_ATTR_NOALIAS);
                        ac_add_attr_dereferenceable(P, UINT64_MAX);
                }
        }
 
        for (i = 0; i < fninfo->num_params; ++i) {
                if (fninfo->assign[i])
                        *fninfo->assign[i] = LLVMGetParam(ctx->main_fn, i);
        }
 
-       si_llvm_add_attribute(ctx->main_fn, "amdgpu-32bit-address-high-bits",
-                             ctx->screen->info.address32_hi);
+       if (ctx->screen->info.address32_hi) {
+               ac_llvm_add_target_dep_function_attr(ctx->main_fn,
+                                                    
"amdgpu-32bit-address-high-bits",
+                                                    
ctx->screen->info.address32_hi);
+       }
 
        if (max_workgroup_size) {
                si_llvm_add_attribute(ctx->main_fn, 
"amdgpu-max-work-group-size",
                                      max_workgroup_size);
        }
        LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
                                           "no-signed-zeros-fp-math",
                                           "true");
 
        if (ctx->screen->debug_flags & DBG(UNSAFE_MATH)) {
-- 
2.7.4

_______________________________________________
mesa-dev mailing list
mesa-dev@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/mesa-dev