diff --git a/src/amd/compiler/aco_ir.cpp b/src/amd/compiler/aco_ir.cpp index d16fd4d648e..a726e8164b3 100644 --- a/src/amd/compiler/aco_ir.cpp +++ b/src/amd/compiler/aco_ir.cpp @@ -7,6 +7,7 @@ #include "aco_ir.h" #include "aco_builder.h" +#include "aco_shader_info.h" #include "util/u_debug.h" @@ -61,42 +62,41 @@ init() void init_program(Program* program, Stage stage, const struct aco_shader_info* info, - enum amd_gfx_level gfx_level, enum radeon_family family, bool wgp_mode, - ac_shader_config* config) + const aco_compiler_options* options, ac_shader_config* config) { - assert(family != CHIP_UNKNOWN); + assert(options->family != CHIP_UNKNOWN); instruction_buffer = &program->m; program->stage = stage; program->config = config; program->info = *info; - program->gfx_level = gfx_level; - program->family = family; + program->gfx_level = options->gfx_level; + program->family = options->family; program->wave_size = info->wave_size; program->lane_mask = program->wave_size == 32 ? s1 : s2; /* GFX6: There is 64KB LDS per CU, but a single workgroup can only use 32KB. */ - program->dev.lds_limit = gfx_level >= GFX7 ? 65536 : 32768; + program->dev.lds_limit = program->gfx_level >= GFX7 ? 65536 : 32768; /* apparently gfx702 also has 16-bank LDS but I can't find a family for that */ - program->dev.has_16bank_lds = family == CHIP_KABINI || family == CHIP_STONEY; + program->dev.has_16bank_lds = program->family == CHIP_KABINI || program->family == CHIP_STONEY; program->dev.vgpr_limit = 256; program->dev.physical_vgprs = 256; program->dev.vgpr_alloc_granule = 4; - if (gfx_level >= GFX10) { + if (program->gfx_level >= GFX10) { program->dev.physical_sgprs = 128 * 20; /* enough for max waves */ program->dev.sgpr_alloc_granule = 128; program->dev.sgpr_limit = 108; /* includes VCC, which can be treated as s[106-107] on GFX10+ */ - if (family == CHIP_NAVI31 || family == CHIP_NAVI32 || family == CHIP_STRIX_HALO || - gfx_level >= GFX12) { + if (program->family == CHIP_NAVI31 || program->family == CHIP_NAVI32 || + program->family == CHIP_STRIX_HALO || program->gfx_level >= GFX12) { program->dev.physical_vgprs = program->wave_size == 32 ? 1536 : 768; program->dev.vgpr_alloc_granule = program->wave_size == 32 ? 24 : 12; } else { program->dev.physical_vgprs = program->wave_size == 32 ? 1024 : 512; - if (gfx_level >= GFX10_3) + if (program->gfx_level >= GFX10_3) program->dev.vgpr_alloc_granule = program->wave_size == 32 ? 16 : 8; else program->dev.vgpr_alloc_granule = program->wave_size == 32 ? 8 : 4; @@ -105,7 +105,7 @@ init_program(Program* program, Stage stage, const struct aco_shader_info* info, program->dev.physical_sgprs = 800; program->dev.sgpr_alloc_granule = 16; program->dev.sgpr_limit = 102; - if (family == CHIP_TONGA || family == CHIP_ICELAND) + if (program->family == CHIP_TONGA || program->family == CHIP_ICELAND) program->dev.sgpr_alloc_granule = 96; /* workaround hardware bug */ } else { program->dev.physical_sgprs = 512; @@ -120,7 +120,7 @@ init_program(Program* program, Stage stage, const struct aco_shader_info* info, program->dev.vgpr_limit = util_round_down_npot(vgpr_limit, program->dev.vgpr_alloc_granule); } - program->dev.scratch_alloc_granule = gfx_level >= GFX11 ? 256 : 1024; + program->dev.scratch_alloc_granule = program->gfx_level >= GFX11 ? 256 : 1024; program->dev.max_waves_per_simd = 10; if (program->gfx_level >= GFX10_3) @@ -201,7 +201,7 @@ init_program(Program* program, Stage stage, const struct aco_shader_info* info, program->dev.max_nsa_vgprs = 0; } - program->wgp_mode = wgp_mode; + program->wgp_mode = options->wgp_mode; program->progress = CompilationProgress::after_isel; diff --git a/src/amd/compiler/aco_ir.h b/src/amd/compiler/aco_ir.h index c10e82c6592..c432d37f230 100644 --- a/src/amd/compiler/aco_ir.h +++ b/src/amd/compiler/aco_ir.h @@ -2387,8 +2387,7 @@ struct ra_test_policy { void init(); void init_program(Program* program, Stage stage, const struct aco_shader_info* info, - enum amd_gfx_level gfx_level, enum radeon_family family, bool wgp_mode, - ac_shader_config* config); + const aco_compiler_options* options, ac_shader_config* config); void select_program(Program* program, unsigned shader_count, struct nir_shader* const* shaders, ac_shader_config* config, const struct aco_compiler_options* options, diff --git a/src/amd/compiler/instruction_selection/aco_isel_setup.cpp b/src/amd/compiler/instruction_selection/aco_isel_setup.cpp index 4d559b15833..e3a222f8340 100644 --- a/src/amd/compiler/instruction_selection/aco_isel_setup.cpp +++ b/src/amd/compiler/instruction_selection/aco_isel_setup.cpp @@ -740,8 +740,7 @@ setup_isel_context(Program* program, unsigned shader_count, struct nir_shader* c } } - init_program(program, Stage{info->hw_stage, sw_stage}, info, options->gfx_level, options->family, - options->wgp_mode, config); + init_program(program, Stage{info->hw_stage, sw_stage}, info, options, config); isel_context ctx = {}; ctx.program = program; diff --git a/src/amd/compiler/instruction_selection/aco_select_rt_prolog.cpp b/src/amd/compiler/instruction_selection/aco_select_rt_prolog.cpp index b930af8e689..ddf24b935e1 100644 --- a/src/amd/compiler/instruction_selection/aco_select_rt_prolog.cpp +++ b/src/amd/compiler/instruction_selection/aco_select_rt_prolog.cpp @@ -16,8 +16,7 @@ select_rt_prolog(Program* program, ac_shader_config* config, const struct aco_compiler_options* options, const struct aco_shader_info* info, const struct ac_shader_args* in_args, const struct ac_shader_args* out_args) { - init_program(program, compute_cs, info, options->gfx_level, options->family, options->wgp_mode, - config); + init_program(program, compute_cs, info, options, config); Block* block = program->create_and_insert_block(); block->kind = block_kind_top_level; program->workgroup_size = info->workgroup_size; diff --git a/src/amd/compiler/instruction_selection/aco_select_trap_handler.cpp b/src/amd/compiler/instruction_selection/aco_select_trap_handler.cpp index e0811c5625c..f2186c7fc7f 100644 --- a/src/amd/compiler/instruction_selection/aco_select_trap_handler.cpp +++ b/src/amd/compiler/instruction_selection/aco_select_trap_handler.cpp @@ -297,8 +297,7 @@ select_trap_handler_shader(Program* program, ac_shader_config* config, assert(options->gfx_level >= GFX8 && options->gfx_level <= GFX12); - init_program(program, compute_cs, info, options->gfx_level, options->family, options->wgp_mode, - config); + init_program(program, compute_cs, info, options, config); isel_context ctx = {}; ctx.program = program; diff --git a/src/amd/compiler/instruction_selection/aco_select_vs_prolog.cpp b/src/amd/compiler/instruction_selection/aco_select_vs_prolog.cpp index 25d763df1d6..f272026f618 100644 --- a/src/amd/compiler/instruction_selection/aco_select_vs_prolog.cpp +++ b/src/amd/compiler/instruction_selection/aco_select_vs_prolog.cpp @@ -366,8 +366,7 @@ select_vs_prolog(Program* program, const struct aco_vs_prolog_info* pinfo, ac_sh /* This should be enough for any shader/stage. */ unsigned max_user_sgprs = options->gfx_level >= GFX9 ? 32 : 16; - init_program(program, compute_cs, info, options->gfx_level, options->family, options->wgp_mode, - config); + init_program(program, compute_cs, info, options, config); program->dev.vgpr_limit = 256; Block* block = program->create_and_insert_block(); diff --git a/src/amd/compiler/tests/helpers.cpp b/src/amd/compiler/tests/helpers.cpp index 0f1c38e41f1..aa0f4d6ddda 100644 --- a/src/amd/compiler/tests/helpers.cpp +++ b/src/amd/compiler/tests/helpers.cpp @@ -13,6 +13,7 @@ #include "drm-shim/amdgpu_noop_drm_shim.h" #include +#include "ac_gpu_info.h" #include #include @@ -93,7 +94,13 @@ create_program(enum amd_gfx_level gfx_level, Stage stage, unsigned wave_size, memset(&config, 0, sizeof(config)); info.wave_size = wave_size; program.reset(new Program); - aco::init_program(program.get(), stage, &info, gfx_level, family, false, &config); + rad_info.gfx_level = gfx_level; + rad_info.family = family; + struct aco_compiler_options options = { + .family = family, + .gfx_level = gfx_level, + }; + aco::init_program(program.get(), stage, &info, &options, &config); program->workgroup_size = UINT_MAX; calc_min_waves(program.get());