| 123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154 |
- // Copyright 2021 yuzu Emulator Project
- // Licensed under GPLv2 or any later version
- // Refer to the license.txt file included.
- #include <algorithm>
- #include <memory>
- #include <vector>
- #include "shader_recompiler/frontend/ir/basic_block.h"
- #include "shader_recompiler/frontend/ir/post_order.h"
- #include "shader_recompiler/frontend/maxwell/program.h"
- #include "shader_recompiler/frontend/maxwell/structured_control_flow.h"
- #include "shader_recompiler/frontend/maxwell/translate/translate.h"
- #include "shader_recompiler/ir_opt/passes.h"
- namespace Shader::Maxwell {
- namespace {
- void RemoveUnreachableBlocks(IR::Program& program) {
- // Some blocks might be unreachable if a function call exists unconditionally
- // If this happens the number of blocks and post order blocks will mismatch
- if (program.blocks.size() == program.post_order_blocks.size()) {
- return;
- }
- const auto begin{program.blocks.begin() + 1};
- const auto end{program.blocks.end()};
- const auto pred{[](IR::Block* block) { return block->ImmediatePredecessors().empty(); }};
- program.blocks.erase(std::remove_if(begin, end, pred), end);
- }
- void CollectInterpolationInfo(Environment& env, IR::Program& program) {
- if (program.stage != Stage::Fragment) {
- return;
- }
- const ProgramHeader& sph{env.SPH()};
- for (size_t index = 0; index < program.info.input_generics.size(); ++index) {
- std::optional<PixelImap> imap;
- for (const PixelImap value : sph.ps.GenericInputMap(static_cast<u32>(index))) {
- if (value == PixelImap::Unused) {
- continue;
- }
- if (imap && imap != value) {
- throw NotImplementedException("Per component interpolation");
- }
- imap = value;
- }
- if (!imap) {
- continue;
- }
- program.info.input_generics[index].interpolation = [&] {
- switch (*imap) {
- case PixelImap::Unused:
- case PixelImap::Perspective:
- return Interpolation::Smooth;
- case PixelImap::Constant:
- return Interpolation::Flat;
- case PixelImap::ScreenLinear:
- return Interpolation::NoPerspective;
- }
- throw NotImplementedException("Unknown interpolation {}", *imap);
- }();
- }
- }
- void AddNVNStorageBuffers(IR::Program& program) {
- if (!program.info.uses_global_memory) {
- return;
- }
- const u32 driver_cbuf{0};
- const u32 descriptor_size{0x10};
- const u32 num_buffers{16};
- const u32 base{[&] {
- switch (program.stage) {
- case Stage::VertexA:
- case Stage::VertexB:
- return 0x110u;
- case Stage::TessellationControl:
- return 0x210u;
- case Stage::TessellationEval:
- return 0x310u;
- case Stage::Geometry:
- return 0x410u;
- case Stage::Fragment:
- return 0x510u;
- case Stage::Compute:
- return 0x310u;
- }
- throw InvalidArgument("Invalid stage {}", program.stage);
- }()};
- auto& descs{program.info.storage_buffers_descriptors};
- for (u32 index = 0; index < num_buffers; ++index) {
- const u32 offset{base + index * descriptor_size};
- const auto it{std::ranges::find(descs, offset, &StorageBufferDescriptor::cbuf_offset)};
- if (it != descs.end()) {
- continue;
- }
- // Assume these are written for now
- descs.push_back({
- .cbuf_index = driver_cbuf,
- .cbuf_offset = offset,
- .count = 1,
- .is_written = true,
- });
- }
- }
- } // Anonymous namespace
- IR::Program TranslateProgram(ObjectPool<IR::Inst>& inst_pool, ObjectPool<IR::Block>& block_pool,
- Environment& env, Flow::CFG& cfg) {
- IR::Program program;
- program.blocks = VisitAST(inst_pool, block_pool, env, cfg);
- program.post_order_blocks = PostOrder(program.blocks);
- program.stage = env.ShaderStage();
- program.local_memory_size = env.LocalMemorySize();
- switch (program.stage) {
- case Stage::TessellationControl: {
- const ProgramHeader& sph{env.SPH()};
- program.invocations = sph.common2.threads_per_input_primitive;
- break;
- }
- case Stage::Geometry: {
- const ProgramHeader& sph{env.SPH()};
- program.output_topology = sph.common3.output_topology;
- program.output_vertices = sph.common4.max_output_vertices;
- program.invocations = sph.common2.threads_per_input_primitive;
- break;
- }
- case Stage::Compute:
- program.workgroup_size = env.WorkgroupSize();
- program.shared_memory_size = env.SharedMemorySize();
- break;
- default:
- break;
- }
- RemoveUnreachableBlocks(program);
- // Replace instructions before the SSA rewrite
- Optimization::LowerFp16ToFp32(program);
- Optimization::SsaRewritePass(program);
- Optimization::GlobalMemoryToStorageBufferPass(program);
- Optimization::TexturePass(env, program);
- Optimization::ConstantPropagationPass(program);
- Optimization::DeadCodeEliminationPass(program);
- Optimization::IdentityRemovalPass(program);
- Optimization::VerificationPass(program);
- Optimization::CollectShaderInfoPass(env, program);
- CollectInterpolationInfo(env, program);
- AddNVNStorageBuffers(program);
- return program;
- }
- } // namespace Shader::Maxwell
|