[Mesa-dev] [PATCH] radv: take LDS into account for compute shader occupancy stats
Samuel Pitoiset
samuel.pitoiset at gmail.com
Fri Feb 1 14:34:46 UTC 2019
Fixed.
On 2/1/19 2:57 PM, Mike Lothian wrote:
> Hi
>
> I think you've left a few references to the old ac_nir_get_max_workgroup_size
>
> FAILED: src/amd/vulkan/9198681@@vulkan_radeon at sha/radv_nir_to_llvm.c.o
> x86_64-pc-linux-gnu-gcc -m32
> -Isrc/amd/vulkan/9198681@@vulkan_radeon at sha -Isrc/amd/vulkan
> -I../mesa-9999/src/amd/vulkan -Isrc/../include
> -I../mesa-9999/src/../include -Isrc -I../mesa-9999/src -Isrc/mapi
> rc/mesa -I../mesa-9999/src/gallium/include -Isrc/gallium/auxiliary
> -I../mesa-9999/src/gallium/auxiliary -Isrc/amd -I../mesa-9999/src/amd
> -Isrc/amd/common -I../mesa-9999/src/amd/common -Isrc/compiler -I.
> -9999/src/vulkan/util -Isrc/vulkan/wsi -I../mesa-9999/src/vulkan/wsi
> -Isrc/compiler/nir -I../mesa-9999/src/compiler/nir
> -I/usr/lib/llvm/9/include -I/usr/include/libdrm
> -fdiagnostics-color=always -DNDEBU
> '-DPACKAGE_BUGREPORT="https://bugs.freedesktop.org/enter_bug.cgi?product=Mesa"'
> -DGLX_USE_TLS -DHAVE_ST_VDPAU -DENABLE_ST_OMX_BELLAGIO=0
> -DENABLE_ST_OMX_TIZONIA=0 -DHAVE_X11_PLATFORM -DGLX_INDIRECT_REND
> M_PLATFORM -DHAVE_SURFACELESS_PLATFORM -DENABLE_SHADER_CACHE
> -DHAVE___BUILTIN_BSWAP32 -DHAVE___BUILTIN_BSWAP64 -DHAVE___BUILTIN_CLZ
> -DHAVE___BUILTIN_CLZLL -DHAVE___BUILTIN_CTZ -DHAVE___BUILTIN_EXPECT -D
> UILTIN_POPCOUNT -DHAVE___BUILTIN_POPCOUNTLL
> -DHAVE___BUILTIN_UNREACHABLE -DHAVE_FUNC_ATTRIBUTE_CONST
> -DHAVE_FUNC_ATTRIBUTE_FLATTEN -DHAVE_FUNC_ATTRIBUTE_MALLOC
> -DHAVE_FUNC_ATTRIBUTE_PURE -DHAVE_FUNC_ATT
> LT -DHAVE_FUNC_ATTRIBUTE_WEAK -DHAVE_FUNC_ATTRIBUTE_FORMAT
> -DHAVE_FUNC_ATTRIBUTE_PACKED -DHAVE_FUNC_ATTRIBUTE_RETURNS_NONNULL
> -DHAVE_FUNC_ATTRIBUTE_VISIBILITY -DHAVE_FUNC_ATTRIBUTE_ALIAS
> -DHAVE_FUNC_ATT
> S -DUSE_X86_ASM -DUSE_MMX_ASM -DUSE_3DNOW_ASM -DUSE_SSE_ASM
> -DMAJOR_IN_SYSMACROS -DHAVE_SYS_SYSCTL_H -DHAVE_LINUX_FUTEX_H
> -DHAVE_ENDIAN_H -DHAVE_DLFCN_H -DHAVE_STRTOF -DHAVE_MKOSTEMP
> -DHAVE_POSIX_MEMALI
> RTOD_L -DHAVE_DLADDR -DHAVE_DL_ITERATE_PHDR -DHAVE_ZLIB -DHAVE_PTHREAD
> -DHAVE_PTHREAD_SETAFFINITY -DHAVE_LIBDRM -DHAVE_LLVM=0x0900
> -DMESA_LLVM_VERSION_PATCH=0 -DHAVE_WAYLAND_PLATFORM
> -DWL_HIDE_DEPRECATE
> S=1 -Werror=implicit-function-declaration -Werror=missing-prototypes
> -Werror=return-type -fno-math-errno -fno-trapping-math
> -Wno-missing-field-initializers -Wno-format-truncation -fPIC -pthread
> -D_FILE_
> -D__STDC_LIMIT_MACROS -D_LARGEFILE_SOURCE -D__STDC_CONSTANT_MACROS
> -fvisibility=hidden -Wno-override-init -DVK_USE_PLATFORM_XCB_KHR
> -DVK_USE_PLATFORM_XLIB_KHR -DVK_USE_PLATFORM_WAYLAND_KHR -DVK_USE_PLAT
> O3 -march=native -pipe -MD -MQ
> 'src/amd/vulkan/9198681@@vulkan_radeon at sha/radv_nir_to_llvm.c.o' -MF
> 'src/amd/vulkan/9198681@@vulkan_radeon at sha/radv_nir_to_llvm.c.o.d' -o
> 'src/amd/vulkan/9198681@@vulkan_
> /amd/vulkan/radv_nir_to_llvm.c
> In file included from ../mesa-9999/src/mesa/main/macros.h:35,
> from ../mesa-9999/src/amd/vulkan/radv_private.h:51,
> from ../mesa-9999/src/amd/vulkan/radv_nir_to_llvm.c:28:
> ../mesa-9999/src/amd/vulkan/radv_nir_to_llvm.c: In function
> ‘ac_translate_nir_to_llvm’:
> ../mesa-9999/src/amd/vulkan/radv_nir_to_llvm.c:3453:33: error:
> implicit declaration of function ‘ac_nir_get_max_workgroup_size’; did
> you mean ‘radv_nir_get_max_workgroup_size’? [-Werror=implicit-functio
>
> ac_nir_get_max_workgroup_size(ctx.options->chip_class,
> ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~
> ../mesa-9999/src/util/u_math.h:659:31: note: in definition of macro ‘MAX2’
> #define MAX2( A, B ) ( (A)>(B) ? (A) : (B) )
> ^
> cc1: some warnings being treated as errors
>
> Cheers
>
> Mike
>
> On Fri, 1 Feb 2019 at 11:07, Timothy Arceri <tarceri at itsqueeze.com> wrote:
>> Ported from d205faeb6c96.
>> ---
>> src/amd/vulkan/radv_nir_to_llvm.c | 6 +++---
>> src/amd/vulkan/radv_private.h | 3 +++
>> src/amd/vulkan/radv_shader.c | 10 ++++++++--
>> 3 files changed, 14 insertions(+), 5 deletions(-)
>>
>> diff --git a/src/amd/vulkan/radv_nir_to_llvm.c b/src/amd/vulkan/radv_nir_to_llvm.c
>> index e80938527e5..d90a4c0de1e 100644
>> --- a/src/amd/vulkan/radv_nir_to_llvm.c
>> +++ b/src/amd/vulkan/radv_nir_to_llvm.c
>> @@ -3372,9 +3372,9 @@ ac_setup_rings(struct radv_shader_context *ctx)
>> }
>> }
>>
>> -static unsigned
>> -ac_nir_get_max_workgroup_size(enum chip_class chip_class,
>> - const struct nir_shader *nir)
>> +unsigned
>> +radv_nir_get_max_workgroup_size(enum chip_class chip_class,
>> + const struct nir_shader *nir)
>> {
>> switch (nir->info.stage) {
>> case MESA_SHADER_TESS_CTRL:
>> diff --git a/src/amd/vulkan/radv_private.h b/src/amd/vulkan/radv_private.h
>> index 85c18906f84..e5b8286ea62 100644
>> --- a/src/amd/vulkan/radv_private.h
>> +++ b/src/amd/vulkan/radv_private.h
>> @@ -1934,6 +1934,9 @@ void radv_compile_nir_shader(struct ac_llvm_compiler *ac_llvm,
>> int nir_count,
>> const struct radv_nir_compiler_options *options);
>>
>> +unsigned radv_nir_get_max_workgroup_size(enum chip_class chip_class,
>> + const struct nir_shader *nir);
>> +
>> /* radv_shader_info.h */
>> struct radv_shader_info;
>>
>> diff --git a/src/amd/vulkan/radv_shader.c b/src/amd/vulkan/radv_shader.c
>> index 07450ff236b..a7fce02ee83 100644
>> --- a/src/amd/vulkan/radv_shader.c
>> +++ b/src/amd/vulkan/radv_shader.c
>> @@ -744,7 +744,8 @@ generate_shader_stats(struct radv_device *device,
>> gl_shader_stage stage,
>> struct _mesa_string_buffer *buf)
>> {
>> - unsigned lds_increment = device->physical_device->rad_info.chip_class >= CIK ? 512 : 256;
>> + enum chip_class chip_class = device->physical_device->rad_info.chip_class;
>> + unsigned lds_increment = chip_class >= CIK ? 512 : 256;
>> struct ac_shader_config *conf;
>> unsigned max_simd_waves;
>> unsigned lds_per_wave = 0;
>> @@ -757,12 +758,17 @@ generate_shader_stats(struct radv_device *device,
>> lds_per_wave = conf->lds_size * lds_increment +
>> align(variant->info.fs.num_interp * 48,
>> lds_increment);
>> + } else if (stage == MESA_SHADER_COMPUTE) {
>> + unsigned max_workgroup_size =
>> + ac_nir_get_max_workgroup_size(chip_class, variant->nir);
>> + lds_per_wave = (conf->lds_size * lds_increment) /
>> + DIV_ROUND_UP(max_workgroup_size, 64);
>> }
>>
>> if (conf->num_sgprs)
>> max_simd_waves =
>> MIN2(max_simd_waves,
>> - ac_get_num_physical_sgprs(device->physical_device->rad_info.chip_class) / conf->num_sgprs);
>> + ac_get_num_physical_sgprs(chip_class) / conf->num_sgprs);
>>
>> if (conf->num_vgprs)
>> max_simd_waves =
>> --
>> 2.20.1
>>
>> _______________________________________________
>> mesa-dev mailing list
>> mesa-dev at lists.freedesktop.org
>> https://lists.freedesktop.org/mailman/listinfo/mesa-dev
> _______________________________________________
> mesa-dev mailing list
> mesa-dev at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/mesa-dev
More information about the mesa-dev
mailing list