851 bool si_shader_binary_upload(struct si_screen *sscreen, struct si_shader *shader,  in si_shader_binary_upload()  argument
855    if (!si_shader_binary_open(sscreen, shader, &binary))  in si_shader_binary_upload()
860       &sscreen->b,  in si_shader_binary_upload()
861       (sscreen->info.cpdma_prefetch_writes_memory ? 0 : SI_RESOURCE_FLAG_READ_ONLY) |  in si_shader_binary_upload()
873    u.rx_ptr = sscreen->ws->buffer_map(sscreen->ws,  in si_shader_binary_upload()
881    if (sscreen->debug_flags & DBG(SQTT)) {  in si_shader_binary_upload()
888    sscreen->ws->buffer_unmap(sscreen->ws, shader->bo->buf);  in si_shader_binary_upload()
956    struct si_screen *sscreen = shader->selector->screen;  in si_calculate_max_simd_waves()  local
959    unsigned lds_increment = sscreen->info.chip_class >= GFX7 ? 512 : 256;  in si_calculate_max_simd_waves()
963    max_simd_waves = sscreen->info.max_wave64_per_simd;  in si_calculate_max_simd_waves()
983                         DIV_ROUND_UP(max_workgroup_size, sscreen->compute_wave_size);  in si_calculate_max_simd_waves()
992          MIN2(max_simd_waves, sscreen->info.num_physical_sgprs_per_simd / conf->num_sgprs);  in si_calculate_max_simd_waves()
998       unsigned max_vgprs = sscreen->info.num_physical_wave64_vgprs_per_simd;  in si_calculate_max_simd_waves()
1002    unsigned max_lds_per_simd = sscreen->info.lds_size_per_workgroup / 4;  in si_calculate_max_simd_waves()
1027 static void si_shader_dump_stats(struct si_screen *sscreen, struct si_shader *shader, FILE *file,  in si_shader_dump_stats()  argument
1032    if (!check_debug_option || si_can_dump_shader(sscreen, shader->selector->info.stage)) {  in si_shader_dump_stats()
1054               shader->info.private_mem_vgprs, si_get_shader_binary_size(sscreen, shader),  in si_shader_dump_stats()
1094 void si_shader_dump(struct si_screen *sscreen, struct si_shader *shader,  in si_shader_dump()  argument
1099    if (!check_debug_option || si_can_dump_shader(sscreen, stage))  in si_shader_dump()
1113        (si_can_dump_shader(sscreen, stage) && !(sscreen->debug_flags & DBG(NO_ASM)))) {  in si_shader_dump()
1119          si_shader_dump_disassembly(sscreen, &shader->prolog->binary, stage, wave_size, debug,  in si_shader_dump()
1122          si_shader_dump_disassembly(sscreen, &shader->previous_stage->binary, stage,  in si_shader_dump()
1125          si_shader_dump_disassembly(sscreen, &shader->prolog2->binary, stage, wave_size,  in si_shader_dump()
1128       si_shader_dump_disassembly(sscreen, &shader->binary, stage, wave_size, debug, "main",  in si_shader_dump()
1132          si_shader_dump_disassembly(sscreen, &shader->epilog->binary, stage, wave_size, debug,  in si_shader_dump()
1137    si_shader_dump_stats(sscreen, shader, file, check_debug_option);  in si_shader_dump()
1411 bool si_compile_shader(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,  in si_compile_shader()  argument
1420    if (si_can_dump_shader(sscreen, sel->info.stage) &&  in si_compile_shader()
1421        !(sscreen->debug_flags & DBG(NO_NIR))) {  in si_compile_shader()
1437    if (!si_llvm_compile_shader(sscreen, compiler, shader, debug, nir, free_nir))  in si_compile_shader()
1479       unsigned wave_size = sscreen->compute_wave_size;  in si_compile_shader()
1481          sscreen->info.num_physical_wave64_vgprs_per_simd * (wave_size == 32 ? 2 : 1);  in si_compile_shader()
1482       unsigned max_sgprs = sscreen->info.num_physical_sgprs_per_simd;  in si_compile_shader()
1518    si_shader_dump_stats_for_shader_db(sscreen, shader, debug);  in si_compile_shader()
1536 si_get_shader_part(struct si_screen *sscreen, struct si_shader_part **list,  in si_get_shader_part()  argument
1544    simple_mtx_lock(&sscreen->shader_parts_mutex);  in si_get_shader_part()
1549          simple_mtx_unlock(&sscreen->shader_parts_mutex);  in si_get_shader_part()
1559    sel.screen = sscreen;  in si_get_shader_part()
1589    si_llvm_context_init(&ctx, sscreen, compiler,  in si_get_shader_part()
1590                         si_get_wave_size(sscreen, stage,  in si_get_shader_part()
1600    if (!si_compile_llvm(sscreen, &result->binary, &result->config, compiler, &ctx.ac, debug,  in si_get_shader_part()
1612    simple_mtx_unlock(&sscreen->shader_parts_mutex);  in si_get_shader_part()
1616 static bool si_get_vs_prolog(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,  in si_get_vs_prolog()  argument
1631       si_get_shader_part(sscreen, &sscreen->vs_prologs, MESA_SHADER_VERTEX, true, &prolog_key,  in si_get_vs_prolog()
1639 static bool si_shader_select_vs_parts(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,  in si_shader_select_vs_parts()  argument
1642    return si_get_vs_prolog(sscreen, compiler, shader, debug, shader, &shader->key.part.vs.prolog);  in si_shader_select_vs_parts()
1648 static bool si_shader_select_tcs_parts(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,  in si_shader_select_tcs_parts()  argument
1651    if (sscreen->info.chip_class >= GFX9) {  in si_shader_select_tcs_parts()
1654       if (!si_get_vs_prolog(sscreen, compiler, shader, debug, ls_main_part,  in si_shader_select_tcs_parts()
1666    shader->epilog = si_get_shader_part(sscreen, &sscreen->tcs_epilogs, MESA_SHADER_TESS_CTRL, false,  in si_shader_select_tcs_parts()
1675 static bool si_shader_select_gs_parts(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,  in si_shader_select_gs_parts()  argument
1678    if (sscreen->info.chip_class >= GFX9) {  in si_shader_select_gs_parts()
1687           !si_get_vs_prolog(sscreen, compiler, shader, debug, es_main_part,  in si_shader_select_gs_parts()
1703       si_get_shader_part(sscreen, &sscreen->gs_prologs, MESA_SHADER_GEOMETRY, true, &prolog_key,  in si_shader_select_gs_parts()
1863 static bool si_shader_select_ps_parts(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,  in si_shader_select_ps_parts()  argument
1875          si_get_shader_part(sscreen, &sscreen->ps_prologs, MESA_SHADER_FRAGMENT, true, &prolog_key,  in si_shader_select_ps_parts()
1885       si_get_shader_part(sscreen, &sscreen->ps_epilogs, MESA_SHADER_FRAGMENT, false, &epilog_key,  in si_shader_select_ps_parts()
1954 void si_multiwave_lds_size_workaround(struct si_screen *sscreen, unsigned *lds_size)  in si_multiwave_lds_size_workaround()  argument
1965    if (sscreen->info.family == CHIP_BONAIRE || sscreen->info.family == CHIP_KABINI)  in si_multiwave_lds_size_workaround()
1969 void si_fix_resource_usage(struct si_screen *sscreen, struct si_shader *shader)  in si_fix_resource_usage()  argument
1976        si_get_max_workgroup_size(shader) > sscreen->compute_wave_size) {  in si_fix_resource_usage()
1977       si_multiwave_lds_size_workaround(sscreen, &shader->config.lds_size);  in si_fix_resource_usage()
1981 bool si_create_shader_variant(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,  in si_create_shader_variant()  argument
2000       if (!si_compile_shader(sscreen, compiler, shader, debug))  in si_create_shader_variant()
2038          if (!si_shader_select_vs_parts(sscreen, compiler, shader, debug))  in si_create_shader_variant()
2042          if (!si_shader_select_tcs_parts(sscreen, compiler, shader, debug))  in si_create_shader_variant()
2048          if (!si_shader_select_gs_parts(sscreen, compiler, shader, debug))  in si_create_shader_variant()
2052          if (!si_shader_select_ps_parts(sscreen, compiler, shader, debug))  in si_create_shader_variant()
2107    } else if (sscreen->info.chip_class >= GFX9 && sel->info.stage == MESA_SHADER_GEOMETRY) {  in si_create_shader_variant()
2112       sscreen->use_ngg &&  in si_create_shader_variant()
2120    shader->uses_vs_state_outprim = sscreen->use_ngg &&  in si_create_shader_variant()
2141    si_fix_resource_usage(sscreen, shader);  in si_create_shader_variant()
2142    si_shader_dump(sscreen, shader, debug, stderr, true);  in si_create_shader_variant()
2145    if (!si_shader_binary_upload(sscreen, shader, 0)) {  in si_create_shader_variant()