aco/tests: add isel test helpers

Having NIR-based instruction selection tests instead of GLSL lets us be
more precise with the input NIR.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28301>

src/amd/compiler/tests/helpers.cpp

@@ -26,6 +26,11 @@ std::unique_ptr<Program> program;
 Builder bld(NULL);
 Temp inputs[16];
 
+static radeon_info rad_info;
+static nir_shader_compiler_options nir_options;
+static nir_builder _nb;
+nir_builder *nb;
+
 static VkInstance instance_cache[CHIP_LAST] = {VK_NULL_HANDLE};
 static VkDevice device_cache[CHIP_LAST] = {VK_NULL_HANDLE};
 static std::mutex create_device_mutex;
@@ -118,6 +123,40 @@ setup_cs(const char* input_spec, enum amd_gfx_level gfx_level, enum radeon_famil
    return true;
 }
 
+bool
+setup_nir_cs(enum amd_gfx_level gfx_level, gl_shader_stage stage, enum radeon_family family, const char* subvariant)
+{
+   if (!set_variant(gfx_level, subvariant))
+      return false;
+
+   if (family == CHIP_UNKNOWN) {
+      switch (gfx_level) {
+      case GFX6: family = CHIP_TAHITI; break;
+      case GFX7: family = CHIP_BONAIRE; break;
+      case GFX8: family = CHIP_POLARIS10; break;
+      case GFX9: family = CHIP_VEGA10; break;
+      case GFX10: family = CHIP_NAVI10; break;
+      case GFX10_3: family = CHIP_NAVI21; break;
+      case GFX11: family = CHIP_NAVI31; break;
+      default: family = CHIP_UNKNOWN; break;
+      }
+   }
+
+   memset(&rad_info, 0, sizeof(rad_info));
+   rad_info.gfx_level = gfx_level;
+   rad_info.family = family;
+
+   memset(&nir_options, 0, sizeof(nir_options));
+   ac_set_nir_options(&rad_info, false, &nir_options);
+
+   glsl_type_singleton_init_or_ref();
+
+   _nb = nir_builder_init_simple_shader(stage, &nir_options, "aco_test");
+   nb = &_nb;
+
+   return true;
+}
+
 void
 finish_program(Program* prog, bool endpgm)
 {
@@ -274,6 +313,62 @@ finish_assembler_test()
    }
 }
 
+void
+live_var_analysis_debug_func(void* private_data, enum aco_compiler_debug_level level, const char* message)
+{
+   if (level == ACO_COMPILER_DEBUG_LEVEL_ERROR)
+      *(bool *)private_data = true;
+}
+
+void
+finish_isel_test(enum ac_hw_stage hw_stage, unsigned wave_size)
+{
+   nir_validate_shader(nb->shader, "in finish_isel_test");
+   nir_validate_ssa_dominance(nb->shader, "in finish_isel_test");
+
+   program.reset(new Program);
+   program->debug.func = nullptr;
+   program->debug.private_data = nullptr;
+
+   ac_shader_args args = {};
+
+   aco_compiler_options options = {};
+   options.family = rad_info.family;
+   options.gfx_level = rad_info.gfx_level;
+
+   memset(&info, 0, sizeof(info));
+   info.hw_stage = hw_stage;
+   info.wave_size = wave_size;
+   info.workgroup_size = nb->shader->info.workgroup_size[0] * nb->shader->info.workgroup_size[1] * nb->shader->info.workgroup_size[2];
+
+   memset(&config, 0, sizeof(config));
+
+   select_program(program.get(), 1, &nb->shader, &config, &options, &info, &args);
+
+   ralloc_free(nb->shader);
+   glsl_type_singleton_decref();
+
+   aco_print_program(program.get(), output);
+
+   if (!aco::validate_ir(program.get())) {
+      fail_test("Validation after instruction selection failed");
+      return;
+   }
+   if (!aco::validate_cfg(program.get())) {
+      fail_test("Invalidate CFG");
+      return;
+   }
+
+   bool live_var_fail = false;
+   program->debug.func = &live_var_analysis_debug_func;
+   program->debug.private_data = &live_var_fail;
+   aco::live_var_analysis(program.get());
+   if (live_var_fail) {
+      fail_test("Live var analysis failed");
+      return;
+   }
+}
+
 void
 writeout(unsigned i, Temp tmp)
 {

src/amd/compiler/tests/helpers.h

@@ -9,6 +9,8 @@
 #include "vulkan/vulkan.h"
 
 #include "framework.h"
+#include "ac_gpu_info.h"
+#include "nir_builder.h"
 #include <functional>
 
 enum QoShaderDeclType {
@@ -49,6 +51,7 @@ extern aco_shader_info info;
 extern std::unique_ptr<aco::Program> program;
 extern aco::Builder bld;
 extern aco::Temp inputs[16];
+extern nir_builder *nb;
 
 namespace aco {
 struct ra_test_policy;
@@ -59,6 +62,9 @@ void create_program(enum amd_gfx_level gfx_level, aco::Stage stage, unsigned wav
 bool setup_cs(const char* input_spec, enum amd_gfx_level gfx_level,
               enum radeon_family family = CHIP_UNKNOWN, const char* subvariant = "",
               unsigned wave_size = 64);
+bool
+setup_nir_cs(enum amd_gfx_level gfx_level, gl_shader_stage stage = MESA_SHADER_COMPUTE,
+             enum radeon_family family = CHIP_UNKNOWN, const char* subvariant = "");
 
 void finish_program(aco::Program* program, bool endpgm = true);
 void finish_validator_test();
@@ -72,6 +78,7 @@ void finish_waitcnt_test();
 void finish_insert_nops_test(bool endpgm = true);
 void finish_form_hard_clause_test();
 void finish_assembler_test();
+void finish_isel_test(enum ac_hw_stage hw_stage = AC_HW_COMPUTE_SHADER, unsigned wave_size = 64);
 
 void writeout(unsigned i, aco::Temp tmp = aco::Temp(0, aco::s1));
 void writeout(unsigned i, aco::Builder::Result res);