[Mesa-dev] [PATCH 14/16] ac: use amdgpu-flat-work-group-size

[Marek Olšák]

Yes good point.

Marek

On Fri, Oct 13, 2017 at 3:39 PM, Ernst Sjöstrand <ernstp at gmail.com> wrote:
> Isn't the logic inverted here?
>
> 2017-10-13 14:04 GMT+02:00 Marek Olšák <maraeo at gmail.com>:
>> From: Marek Olšák <marek.olsak at amd.com>
>>
>> the old one is being deprecated or removed
>> ---
>>  src/amd/common/ac_nir_to_llvm.c          | 3 ++-
>>  src/gallium/drivers/radeonsi/si_shader.c | 4 +++-
>>  2 files changed, 5 insertions(+), 2 deletions(-)
>>
>> diff --git a/src/amd/common/ac_nir_to_llvm.c b/src/amd/common/ac_nir_to_llvm.c
>> index 11ba487..4492d8e 100644
>> --- a/src/amd/common/ac_nir_to_llvm.c
>> +++ b/src/amd/common/ac_nir_to_llvm.c
>> @@ -346,21 +346,22 @@ create_llvm_function(LLVMContextRef ctx, LLVMModuleRef module,
>>                         ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_BYVAL);
>>                         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) {
>>                 ac_llvm_add_target_dep_function_attr(main_function,
>> -                                                    "amdgpu-max-work-group-size",
>> +                                                    HAVE_LLVM >= 0x0400 ? "amdgpu-max-work-group-size" :
>> +                                                                          "amdgpu-flat-work-group-size",
>>                                                      max_workgroup_size);
>>         }
>>         if (unsafe_math) {
>>                 /* These were copied from some LLVM test. */
>>                 LLVMAddTargetDependentFunctionAttr(main_function,
>>                                                    "less-precise-fpmad",
>>                                                    "true");
>>                 LLVMAddTargetDependentFunctionAttr(main_function,
>>                                                    "no-infs-fp-math",
>>                                                    "true");
>> diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c
>> index ff372ae..506da6f 100644
>> --- a/src/gallium/drivers/radeonsi/si_shader.c
>> +++ b/src/gallium/drivers/radeonsi/si_shader.c
>> @@ -4155,21 +4155,23 @@ static void si_create_function(struct si_shader_context *ctx,
>>                 } else
>>                         lp_add_function_attr(ctx->main_fn, i + 1, LP_FUNC_ATTR_INREG);
>>         }
>>
>>         for (i = 0; i < fninfo->num_params; ++i) {
>>                 if (fninfo->assign[i])
>>                         *fninfo->assign[i] = LLVMGetParam(ctx->main_fn, i);
>>         }
>>
>>         if (max_workgroup_size) {
>> -               si_llvm_add_attribute(ctx->main_fn, "amdgpu-max-work-group-size",
>> +               si_llvm_add_attribute(ctx->main_fn,
>> +                                     HAVE_LLVM >= 0x0400 ? "amdgpu-max-work-group-size" :
>> +                                                           "amdgpu-flat-work-group-size",
>>                                       max_workgroup_size);
>>         }
>>         LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
>>                                            "no-signed-zeros-fp-math",
>>                                            "true");
>>
>>         if (ctx->screen->b.debug_flags & DBG(UNSAFE_MATH)) {
>>                 /* These were copied from some LLVM test. */
>>                 LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
>>                                                    "less-precise-fpmad",
>> --
>> 2.7.4
>>
>> _______________________________________________
>> mesa-dev mailing list
>> mesa-dev at lists.freedesktop.org
>> https://lists.freedesktop.org/mailman/listinfo/mesa-dev