Module: Mesa Branch: master Commit: 486bc1e17ecc975f98fb495bd2f8ae580eebbf6e URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=486bc1e17ecc975f98fb495bd2f8ae580eebbf6e
Author: Marek Olšák <[email protected]> Date: Fri May 31 15:38:39 2019 -0400 ac: use amdgpu-flat-work-group-size Reviewed-by: Bas Nieuwenhuizen <[email protected]> --- src/amd/common/ac_llvm_util.c | 10 ++++++++++ src/amd/common/ac_llvm_util.h | 1 + src/amd/vulkan/radv_nir_to_llvm.c | 7 ++----- src/gallium/drivers/radeonsi/si_shader.c | 7 ++----- 4 files changed, 15 insertions(+), 10 deletions(-) diff --git a/src/amd/common/ac_llvm_util.c b/src/amd/common/ac_llvm_util.c index 5b701603ebb..c8a8bf146fe 100644 --- a/src/amd/common/ac_llvm_util.c +++ b/src/amd/common/ac_llvm_util.c @@ -269,6 +269,16 @@ ac_llvm_add_target_dep_function_attr(LLVMValueRef F, LLVMAddTargetDependentFunctionAttr(F, name, str); } +void ac_llvm_set_workgroup_size(LLVMValueRef F, unsigned size) +{ + if (!size) + return; + + char str[32]; + snprintf(str, sizeof(str), "%u,%u", size, size); + LLVMAddTargetDependentFunctionAttr(F, "amdgpu-flat-work-group-size", str); +} + unsigned ac_count_scratch_private_memory(LLVMValueRef function) { diff --git a/src/amd/common/ac_llvm_util.h b/src/amd/common/ac_llvm_util.h index ca00540da80..18102be5207 100644 --- a/src/amd/common/ac_llvm_util.h +++ b/src/amd/common/ac_llvm_util.h @@ -109,6 +109,7 @@ LLVMBuilderRef ac_create_builder(LLVMContextRef ctx, void ac_llvm_add_target_dep_function_attr(LLVMValueRef F, const char *name, unsigned value); +void ac_llvm_set_workgroup_size(LLVMValueRef F, unsigned size); static inline unsigned ac_get_load_intr_attribs(bool can_speculate) diff --git a/src/amd/vulkan/radv_nir_to_llvm.c b/src/amd/vulkan/radv_nir_to_llvm.c index dca4bebcdd1..98a04e27b25 100644 --- a/src/amd/vulkan/radv_nir_to_llvm.c +++ b/src/amd/vulkan/radv_nir_to_llvm.c @@ -518,11 +518,8 @@ create_llvm_function(LLVMContextRef ctx, LLVMModuleRef module, options->address32_hi); } - if (max_workgroup_size) { - ac_llvm_add_target_dep_function_attr(main_function, - "amdgpu-max-work-group-size", - max_workgroup_size); - } + ac_llvm_set_workgroup_size(main_function, max_workgroup_size); + if (options->unsafe_math) { /* These were copied from some LLVM test. */ LLVMAddTargetDependentFunctionAttr(main_function, diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c index 5bd65e0f65c..e044e180778 100644 --- a/src/gallium/drivers/radeonsi/si_shader.c +++ b/src/gallium/drivers/radeonsi/si_shader.c @@ -4307,11 +4307,8 @@ void si_create_function(struct si_shader_context *ctx, ctx->screen->info.address32_hi); } - if (max_workgroup_size) { - ac_llvm_add_target_dep_function_attr(ctx->main_fn, - "amdgpu-max-work-group-size", - max_workgroup_size); - } + ac_llvm_set_workgroup_size(ctx->main_fn, max_workgroup_size); + LLVMAddTargetDependentFunctionAttr(ctx->main_fn, "no-signed-zeros-fp-math", "true"); _______________________________________________ mesa-commit mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/mesa-commit
