aco: pass aco_compiler_options to init_program()

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38701>

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;
 

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,

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;

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;

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;

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();

src/amd/compiler/tests/helpers.cpp

@@ -13,6 +13,7 @@
 #include "drm-shim/amdgpu_noop_drm_shim.h"
 #include <llvm-c/Target.h>
 
+#include "ac_gpu_info.h"
 #include <mutex>
 #include <stdio.h>
 
@@ -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());