forked from ShuriZma/suyu
1
0
Fork 0

spirv: Remove dependencies on Environment when generating SPIR-V

This commit is contained in:
ReinUsesLisp 2021-03-27 03:08:31 -03:00 committed by ameerj
parent cb6039ccea
commit 675a82416d
5 changed files with 15 additions and 16 deletions

View File

@ -126,12 +126,12 @@ Id DefineMain(EmitContext& ctx, IR::Program& program) {
return main; return main;
} }
void DefineEntryPoint(Environment& env, const IR::Program& program, EmitContext& ctx, Id main) { void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) {
const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size()); const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size());
spv::ExecutionModel execution_model{}; spv::ExecutionModel execution_model{};
switch (program.stage) { switch (program.stage) {
case Shader::Stage::Compute: { case Shader::Stage::Compute: {
const std::array<u32, 3> workgroup_size{env.WorkgroupSize()}; const std::array<u32, 3> workgroup_size{program.workgroup_size};
execution_model = spv::ExecutionModel::GLCompute; execution_model = spv::ExecutionModel::GLCompute;
ctx.AddExecutionMode(main, spv::ExecutionMode::LocalSize, workgroup_size[0], ctx.AddExecutionMode(main, spv::ExecutionMode::LocalSize, workgroup_size[0],
workgroup_size[1], workgroup_size[2]); workgroup_size[1], workgroup_size[2]);
@ -148,7 +148,7 @@ void DefineEntryPoint(Environment& env, const IR::Program& program, EmitContext&
} }
break; break;
default: default:
throw NotImplementedException("Stage {}", env.ShaderStage()); throw NotImplementedException("Stage {}", program.stage);
} }
ctx.AddEntryPoint(execution_model, main, "main", interfaces); ctx.AddEntryPoint(execution_model, main, "main", interfaces);
} }
@ -267,11 +267,10 @@ Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) {
} }
} // Anonymous namespace } // Anonymous namespace
std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program, std::vector<u32> EmitSPIRV(const Profile& profile, IR::Program& program, u32& binding) {
u32& binding) {
EmitContext ctx{profile, program, binding}; EmitContext ctx{profile, program, binding};
const Id main{DefineMain(ctx, program)}; const Id main{DefineMain(ctx, program)};
DefineEntryPoint(env, program, ctx, main); DefineEntryPoint(program, ctx, main);
if (profile.support_float_controls) { if (profile.support_float_controls) {
ctx.AddExtension("SPV_KHR_float_controls"); ctx.AddExtension("SPV_KHR_float_controls");
SetupDenormControl(profile, program, ctx, main); SetupDenormControl(profile, program, ctx, main);

View File

@ -8,15 +8,14 @@
#include "common/common_types.h" #include "common/common_types.h"
#include "shader_recompiler/backend/spirv/emit_context.h" #include "shader_recompiler/backend/spirv/emit_context.h"
#include "shader_recompiler/environment.h"
#include "shader_recompiler/frontend/ir/microinstruction.h" #include "shader_recompiler/frontend/ir/microinstruction.h"
#include "shader_recompiler/frontend/ir/program.h" #include "shader_recompiler/frontend/ir/program.h"
#include "shader_recompiler/profile.h" #include "shader_recompiler/profile.h"
namespace Shader::Backend::SPIRV { namespace Shader::Backend::SPIRV {
[[nodiscard]] std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env, [[nodiscard]] std::vector<u32> EmitSPIRV(const Profile& profile, IR::Program& program,
IR::Program& program, u32& binding); u32& binding);
// Microinstruction emitters // Microinstruction emitters
Id EmitPhi(EmitContext& ctx, IR::Inst* inst); Id EmitPhi(EmitContext& ctx, IR::Inst* inst);

View File

@ -4,6 +4,7 @@
#pragma once #pragma once
#include <array>
#include <string> #include <string>
#include <boost/container/small_vector.hpp> #include <boost/container/small_vector.hpp>
@ -19,6 +20,7 @@ struct Program {
BlockList post_order_blocks; BlockList post_order_blocks;
Info info; Info info;
Stage stage{}; Stage stage{};
std::array<u32, 3> workgroup_size{};
}; };
[[nodiscard]] std::string DumpProgram(const Program& program); [[nodiscard]] std::string DumpProgram(const Program& program);

View File

@ -33,6 +33,9 @@ IR::Program TranslateProgram(ObjectPool<IR::Inst>& inst_pool, ObjectPool<IR::Blo
program.blocks = VisitAST(inst_pool, block_pool, env, cfg); program.blocks = VisitAST(inst_pool, block_pool, env, cfg);
program.post_order_blocks = PostOrder(program.blocks); program.post_order_blocks = PostOrder(program.blocks);
program.stage = env.ShaderStage(); program.stage = env.ShaderStage();
if (program.stage == Stage::Compute) {
program.workgroup_size = env.WorkgroupSize();
}
RemoveUnreachableBlocks(program); RemoveUnreachableBlocks(program);
// Replace instructions before the SSA rewrite // Replace instructions before the SSA rewrite

View File

@ -680,7 +680,6 @@ GraphicsPipeline PipelineCache::CreateGraphicsPipeline(ShaderPools& pools,
std::array<vk::ShaderModule, Maxwell::MaxShaderStage> modules; std::array<vk::ShaderModule, Maxwell::MaxShaderStage> modules;
u32 binding{0}; u32 binding{0};
env_index = 0;
for (size_t index = 0; index < Maxwell::MaxShaderProgram; ++index) { for (size_t index = 0; index < Maxwell::MaxShaderProgram; ++index) {
if (key.unique_hashes[index] == u128{}) { if (key.unique_hashes[index] == u128{}) {
continue; continue;
@ -691,11 +690,8 @@ GraphicsPipeline PipelineCache::CreateGraphicsPipeline(ShaderPools& pools,
const size_t stage_index{index - 1}; const size_t stage_index{index - 1};
infos[stage_index] = &program.info; infos[stage_index] = &program.info;
Shader::Environment& env{*envs[env_index]}; const Shader::Profile profile{MakeProfile(key, program.stage)};
++env_index; const std::vector<u32> code{EmitSPIRV(profile, program, binding)};
const Shader::Profile profile{MakeProfile(key, env.ShaderStage())};
const std::vector<u32> code{EmitSPIRV(profile, env, program, binding)};
modules[stage_index] = BuildShader(device, code); modules[stage_index] = BuildShader(device, code);
} }
return GraphicsPipeline(maxwell3d, gpu_memory, scheduler, buffer_cache, texture_cache, device, return GraphicsPipeline(maxwell3d, gpu_memory, scheduler, buffer_cache, texture_cache, device,
@ -753,7 +749,7 @@ ComputePipeline PipelineCache::CreateComputePipeline(ShaderPools& pools,
Shader::Maxwell::Flow::CFG cfg{env, pools.flow_block, env.StartAddress()}; Shader::Maxwell::Flow::CFG cfg{env, pools.flow_block, env.StartAddress()};
Shader::IR::Program program{TranslateProgram(pools.inst, pools.block, env, cfg)}; Shader::IR::Program program{TranslateProgram(pools.inst, pools.block, env, cfg)};
u32 binding{0}; u32 binding{0};
std::vector<u32> code{EmitSPIRV(base_profile, env, program, binding)}; std::vector<u32> code{EmitSPIRV(base_profile, program, binding)};
return ComputePipeline{device, descriptor_pool, update_descriptor_queue, program.info, return ComputePipeline{device, descriptor_pool, update_descriptor_queue, program.info,
BuildShader(device, code)}; BuildShader(device, code)};
} }