diff --git a/src/amd/compiler/tests/helpers.cpp b/src/amd/compiler/tests/helpers.cpp index d72f31e28a8..4e1d57d231a 100644 --- a/src/amd/compiler/tests/helpers.cpp +++ b/src/amd/compiler/tests/helpers.cpp @@ -26,6 +26,11 @@ std::unique_ptr 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) { diff --git a/src/amd/compiler/tests/helpers.h b/src/amd/compiler/tests/helpers.h index a85b3a214b6..dc37de7b428 100644 --- a/src/amd/compiler/tests/helpers.h +++ b/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 enum QoShaderDeclType { @@ -49,6 +51,7 @@ extern aco_shader_info info; extern std::unique_ptr 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);