Reviewed-by: Bas Nieuwenhuizen <[email protected]> Thanks!
On Fri, May 31, 2019 at 10:08 PM Marek Olšák <[email protected]> wrote: > > From: Marek Olšák <[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 > @@ -262,20 +262,30 @@ ac_dump_module(LLVMModuleRef module) > void > ac_llvm_add_target_dep_function_attr(LLVMValueRef F, > const char *name, unsigned value) > { > char str[16]; > > snprintf(str, sizeof(str), "0x%x", value); > 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) > { > unsigned private_mem_vgprs = 0; > > /* Process all LLVM instructions. */ > LLVMBasicBlockRef bb = LLVMGetFirstBasicBlock(function); > while (bb) { > LLVMValueRef next = LLVMGetFirstInstruction(bb); > > 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 > @@ -102,20 +102,21 @@ void ac_dump_module(LLVMModuleRef module); > LLVMValueRef ac_llvm_get_called_value(LLVMValueRef call); > bool ac_llvm_is_function(LLVMValueRef v); > LLVMModuleRef ac_create_module(LLVMTargetMachineRef tm, LLVMContextRef ctx); > > 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, unsigned value); > +void ac_llvm_set_workgroup_size(LLVMValueRef F, unsigned size); > > 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/amd/vulkan/radv_nir_to_llvm.c > b/src/amd/vulkan/radv_nir_to_llvm.c > index 341f6388f32..6f102647ba8 100644 > --- a/src/amd/vulkan/radv_nir_to_llvm.c > +++ b/src/amd/vulkan/radv_nir_to_llvm.c > @@ -511,25 +511,22 @@ create_llvm_function(LLVMContextRef ctx, LLVMModuleRef > module, > ac_add_attr_dereferenceable(P, UINT64_MAX); > } > } > > if (options->address32_hi) { > ac_llvm_add_target_dep_function_attr(main_function, > > "amdgpu-32bit-address-high-bits", > 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, > "less-precise-fpmad", > "true"); > LLVMAddTargetDependentFunctionAttr(main_function, > "no-infs-fp-math", > "true"); > LLVMAddTargetDependentFunctionAttr(main_function, > "no-nans-fp-math", > diff --git a/src/gallium/drivers/radeonsi/si_shader.c > b/src/gallium/drivers/radeonsi/si_shader.c > index d2927d0254b..1ba6b8b6033 100644 > --- a/src/gallium/drivers/radeonsi/si_shader.c > +++ b/src/gallium/drivers/radeonsi/si_shader.c > @@ -4276,25 +4276,22 @@ void si_create_function(struct si_shader_context *ctx, > if (fninfo->assign[i]) > *fninfo->assign[i] = LLVMGetParam(ctx->main_fn, i); > } > > 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) { > - 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"); > > if (ctx->screen->debug_flags & DBG(UNSAFE_MATH)) { > /* These were copied from some LLVM test. */ > LLVMAddTargetDependentFunctionAttr(ctx->main_fn, > "less-precise-fpmad", > "true"); > LLVMAddTargetDependentFunctionAttr(ctx->main_fn, > -- > 2.17.1 > > _______________________________________________ > mesa-dev mailing list > [email protected] > https://lists.freedesktop.org/mailman/listinfo/mesa-dev _______________________________________________ mesa-dev mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/mesa-dev
