871 bool si_shader_binary_upload(struct si_screen *sscreen, struct si_shader *shader,  in si_shader_binary_upload()  argument
875    if (!si_shader_binary_open(sscreen, shader, &binary))  in si_shader_binary_upload()
880       &sscreen->b,  in si_shader_binary_upload()
881       (sscreen->info.cpdma_prefetch_writes_memory ?  in si_shader_binary_upload()
893    u.rx_ptr = sscreen->ws->buffer_map(  in si_shader_binary_upload()
901    sscreen->ws->buffer_unmap(shader->bo->buf);  in si_shader_binary_upload()
969    struct si_screen *sscreen = shader->selector->screen;  in si_calculate_max_simd_waves()  local
972    unsigned lds_increment = sscreen->info.chip_class >= GFX7 ? 512 : 256;  in si_calculate_max_simd_waves()
976    max_simd_waves = sscreen->info.max_wave64_per_simd;  in si_calculate_max_simd_waves()
996                         DIV_ROUND_UP(max_workgroup_size, sscreen->compute_wave_size);  in si_calculate_max_simd_waves()
1005          MIN2(max_simd_waves, sscreen->info.num_physical_sgprs_per_simd / conf->num_sgprs);  in si_calculate_max_simd_waves()
1011       unsigned max_vgprs = sscreen->info.num_physical_wave64_vgprs_per_simd;  in si_calculate_max_simd_waves()
1015    unsigned max_lds_per_simd = sscreen->info.lds_size_per_workgroup / 4;  in si_calculate_max_simd_waves()
1040 static void si_shader_dump_stats(struct si_screen *sscreen, struct si_shader *shader, FILE *file,  in si_shader_dump_stats()  argument
1045    if (!check_debug_option || si_can_dump_shader(sscreen, shader->selector->info.stage)) {  in si_shader_dump_stats()
1067               shader->info.private_mem_vgprs, si_get_shader_binary_size(sscreen, shader),  in si_shader_dump_stats()
1109 void si_shader_dump(struct si_screen *sscreen, struct si_shader *shader,  in si_shader_dump()  argument
1114    if (!check_debug_option || si_can_dump_shader(sscreen, stage))  in si_shader_dump()
1128        (si_can_dump_shader(sscreen, stage) && !(sscreen->debug_flags & DBG(NO_ASM)))) {  in si_shader_dump()
1134          si_shader_dump_disassembly(sscreen, &shader->prolog->binary, stage, wave_size, debug,  in si_shader_dump()
1137          si_shader_dump_disassembly(sscreen, &shader->previous_stage->binary, stage,  in si_shader_dump()
1140          si_shader_dump_disassembly(sscreen, &shader->prolog2->binary, stage, wave_size,  in si_shader_dump()
1143       si_shader_dump_disassembly(sscreen, &shader->binary, stage, wave_size, debug, "main",  in si_shader_dump()
1147          si_shader_dump_disassembly(sscreen, &shader->epilog->binary, stage, wave_size, debug,  in si_shader_dump()
1152    si_shader_dump_stats(sscreen, shader, file, check_debug_option);  in si_shader_dump()
1672 static bool si_llvm_compile_shader(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,  in si_llvm_compile_shader()  argument
1679    si_llvm_context_init(&ctx, sscreen, compiler, si_get_shader_wave_size(shader));  in si_llvm_compile_shader()
1738       if (sscreen->info.chip_class >= GFX9) {  in si_llvm_compile_shader()
1882    if ((debug && debug->debug_message) || si_can_dump_shader(sscreen, ctx.stage)) {  in si_llvm_compile_shader()
1890    if (!si_compile_llvm(sscreen, &shader->binary, &shader->config, compiler, &ctx.ac, debug,  in si_llvm_compile_shader()
1902 bool si_compile_shader(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,  in si_compile_shader()  argument
1911    if (si_can_dump_shader(sscreen, sel->info.stage) &&  in si_compile_shader()
1912        !(sscreen->debug_flags & DBG(NO_NIR))) {  in si_compile_shader()
1926    if (!si_llvm_compile_shader(sscreen, compiler, shader, debug, nir, free_nir))  in si_compile_shader()
1933       unsigned wave_size = sscreen->compute_wave_size;  in si_compile_shader()
1935          sscreen->info.num_physical_wave64_vgprs_per_simd * (wave_size == 32 ? 2 : 1);  in si_compile_shader()
1936       unsigned max_sgprs = sscreen->info.num_physical_sgprs_per_simd;  in si_compile_shader()
1972    si_shader_dump_stats_for_shader_db(sscreen, shader, debug);  in si_compile_shader()
1990 si_get_shader_part(struct si_screen *sscreen, struct si_shader_part **list,  in si_get_shader_part()  argument
1998    simple_mtx_lock(&sscreen->shader_parts_mutex);  in si_get_shader_part()
2003          simple_mtx_unlock(&sscreen->shader_parts_mutex);  in si_get_shader_part()
2013    sel.screen = sscreen;  in si_get_shader_part()
2047    si_llvm_context_init(&ctx, sscreen, compiler,  in si_get_shader_part()
2048                         si_get_wave_size(sscreen, stage,  in si_get_shader_part()
2060    if (!si_compile_llvm(sscreen, &result->binary, &result->config, compiler, &ctx.ac, debug,  in si_get_shader_part()
2072    simple_mtx_unlock(&sscreen->shader_parts_mutex);  in si_get_shader_part()
2076 static bool si_get_vs_prolog(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,  in si_get_vs_prolog()  argument
2091       si_get_shader_part(sscreen, &sscreen->vs_prologs, MESA_SHADER_VERTEX, true, &prolog_key,  in si_get_vs_prolog()
2099 static bool si_shader_select_vs_parts(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,  in si_shader_select_vs_parts()  argument
2102    return si_get_vs_prolog(sscreen, compiler, shader, debug, shader, &shader->key.part.vs.prolog);  in si_shader_select_vs_parts()
2108 static bool si_shader_select_tcs_parts(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,  in si_shader_select_tcs_parts()  argument
2111    if (sscreen->info.chip_class >= GFX9) {  in si_shader_select_tcs_parts()
2114       if (!si_get_vs_prolog(sscreen, compiler, shader, debug, ls_main_part,  in si_shader_select_tcs_parts()
2126    shader->epilog = si_get_shader_part(sscreen, &sscreen->tcs_epilogs, MESA_SHADER_TESS_CTRL, false,  in si_shader_select_tcs_parts()
2135 static bool si_shader_select_gs_parts(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,  in si_shader_select_gs_parts()  argument
2138    if (sscreen->info.chip_class >= GFX9) {  in si_shader_select_gs_parts()
2147           !si_get_vs_prolog(sscreen, compiler, shader, debug, es_main_part,  in si_shader_select_gs_parts()
2163       si_get_shader_part(sscreen, &sscreen->gs_prologs, MESA_SHADER_GEOMETRY, true, &prolog_key,  in si_shader_select_gs_parts()
2323 static bool si_shader_select_ps_parts(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,  in si_shader_select_ps_parts()  argument
2335          si_get_shader_part(sscreen, &sscreen->ps_prologs, MESA_SHADER_FRAGMENT, true, &prolog_key,  in si_shader_select_ps_parts()
2345       si_get_shader_part(sscreen, &sscreen->ps_epilogs, MESA_SHADER_FRAGMENT, false, &epilog_key,  in si_shader_select_ps_parts()
2414 void si_multiwave_lds_size_workaround(struct si_screen *sscreen, unsigned *lds_size)  in si_multiwave_lds_size_workaround()  argument
2425    if (sscreen->info.family == CHIP_BONAIRE || sscreen->info.family == CHIP_KABINI)  in si_multiwave_lds_size_workaround()
2429 void si_fix_resource_usage(struct si_screen *sscreen, struct si_shader *shader)  in si_fix_resource_usage()  argument
2436        si_get_max_workgroup_size(shader) > sscreen->compute_wave_size) {  in si_fix_resource_usage()
2437       si_multiwave_lds_size_workaround(sscreen, &shader->config.lds_size);  in si_fix_resource_usage()
2441 bool si_create_shader_variant(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,  in si_create_shader_variant()  argument
2460       if (!si_compile_shader(sscreen, compiler, shader, debug))  in si_create_shader_variant()
2498          if (!si_shader_select_vs_parts(sscreen, compiler, shader, debug))  in si_create_shader_variant()
2502          if (!si_shader_select_tcs_parts(sscreen, compiler, shader, debug))  in si_create_shader_variant()
2508          if (!si_shader_select_gs_parts(sscreen, compiler, shader, debug))  in si_create_shader_variant()
2512          if (!si_shader_select_ps_parts(sscreen, compiler, shader, debug))  in si_create_shader_variant()
2567    } else if (sscreen->info.chip_class >= GFX9 && sel->info.stage == MESA_SHADER_GEOMETRY) {  in si_create_shader_variant()
2571    si_fix_resource_usage(sscreen, shader);  in si_create_shader_variant()
2572    si_shader_dump(sscreen, shader, debug, stderr, true);  in si_create_shader_variant()
2575    if (!si_shader_binary_upload(sscreen, shader, 0)) {  in si_create_shader_variant()