program.cpp 5.3 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154
  1. // Copyright 2021 yuzu Emulator Project
  2. // Licensed under GPLv2 or any later version
  3. // Refer to the license.txt file included.
  4. #include <algorithm>
  5. #include <memory>
  6. #include <vector>
  7. #include "shader_recompiler/frontend/ir/basic_block.h"
  8. #include "shader_recompiler/frontend/ir/post_order.h"
  9. #include "shader_recompiler/frontend/maxwell/program.h"
  10. #include "shader_recompiler/frontend/maxwell/structured_control_flow.h"
  11. #include "shader_recompiler/frontend/maxwell/translate/translate.h"
  12. #include "shader_recompiler/ir_opt/passes.h"
  13. namespace Shader::Maxwell {
  14. namespace {
  15. void RemoveUnreachableBlocks(IR::Program& program) {
  16. // Some blocks might be unreachable if a function call exists unconditionally
  17. // If this happens the number of blocks and post order blocks will mismatch
  18. if (program.blocks.size() == program.post_order_blocks.size()) {
  19. return;
  20. }
  21. const auto begin{program.blocks.begin() + 1};
  22. const auto end{program.blocks.end()};
  23. const auto pred{[](IR::Block* block) { return block->ImmediatePredecessors().empty(); }};
  24. program.blocks.erase(std::remove_if(begin, end, pred), end);
  25. }
  26. void CollectInterpolationInfo(Environment& env, IR::Program& program) {
  27. if (program.stage != Stage::Fragment) {
  28. return;
  29. }
  30. const ProgramHeader& sph{env.SPH()};
  31. for (size_t index = 0; index < program.info.input_generics.size(); ++index) {
  32. std::optional<PixelImap> imap;
  33. for (const PixelImap value : sph.ps.GenericInputMap(static_cast<u32>(index))) {
  34. if (value == PixelImap::Unused) {
  35. continue;
  36. }
  37. if (imap && imap != value) {
  38. throw NotImplementedException("Per component interpolation");
  39. }
  40. imap = value;
  41. }
  42. if (!imap) {
  43. continue;
  44. }
  45. program.info.input_generics[index].interpolation = [&] {
  46. switch (*imap) {
  47. case PixelImap::Unused:
  48. case PixelImap::Perspective:
  49. return Interpolation::Smooth;
  50. case PixelImap::Constant:
  51. return Interpolation::Flat;
  52. case PixelImap::ScreenLinear:
  53. return Interpolation::NoPerspective;
  54. }
  55. throw NotImplementedException("Unknown interpolation {}", *imap);
  56. }();
  57. }
  58. }
  59. void AddNVNStorageBuffers(IR::Program& program) {
  60. if (!program.info.uses_global_memory) {
  61. return;
  62. }
  63. const u32 driver_cbuf{0};
  64. const u32 descriptor_size{0x10};
  65. const u32 num_buffers{16};
  66. const u32 base{[&] {
  67. switch (program.stage) {
  68. case Stage::VertexA:
  69. case Stage::VertexB:
  70. return 0x110u;
  71. case Stage::TessellationControl:
  72. return 0x210u;
  73. case Stage::TessellationEval:
  74. return 0x310u;
  75. case Stage::Geometry:
  76. return 0x410u;
  77. case Stage::Fragment:
  78. return 0x510u;
  79. case Stage::Compute:
  80. return 0x310u;
  81. }
  82. throw InvalidArgument("Invalid stage {}", program.stage);
  83. }()};
  84. auto& descs{program.info.storage_buffers_descriptors};
  85. for (u32 index = 0; index < num_buffers; ++index) {
  86. const u32 offset{base + index * descriptor_size};
  87. const auto it{std::ranges::find(descs, offset, &StorageBufferDescriptor::cbuf_offset)};
  88. if (it != descs.end()) {
  89. continue;
  90. }
  91. // Assume these are written for now
  92. descs.push_back({
  93. .cbuf_index = driver_cbuf,
  94. .cbuf_offset = offset,
  95. .count = 1,
  96. .is_written = true,
  97. });
  98. }
  99. }
  100. } // Anonymous namespace
  101. IR::Program TranslateProgram(ObjectPool<IR::Inst>& inst_pool, ObjectPool<IR::Block>& block_pool,
  102. Environment& env, Flow::CFG& cfg) {
  103. IR::Program program;
  104. program.blocks = VisitAST(inst_pool, block_pool, env, cfg);
  105. program.post_order_blocks = PostOrder(program.blocks);
  106. program.stage = env.ShaderStage();
  107. program.local_memory_size = env.LocalMemorySize();
  108. switch (program.stage) {
  109. case Stage::TessellationControl: {
  110. const ProgramHeader& sph{env.SPH()};
  111. program.invocations = sph.common2.threads_per_input_primitive;
  112. break;
  113. }
  114. case Stage::Geometry: {
  115. const ProgramHeader& sph{env.SPH()};
  116. program.output_topology = sph.common3.output_topology;
  117. program.output_vertices = sph.common4.max_output_vertices;
  118. program.invocations = sph.common2.threads_per_input_primitive;
  119. break;
  120. }
  121. case Stage::Compute:
  122. program.workgroup_size = env.WorkgroupSize();
  123. program.shared_memory_size = env.SharedMemorySize();
  124. break;
  125. default:
  126. break;
  127. }
  128. RemoveUnreachableBlocks(program);
  129. // Replace instructions before the SSA rewrite
  130. Optimization::LowerFp16ToFp32(program);
  131. Optimization::SsaRewritePass(program);
  132. Optimization::GlobalMemoryToStorageBufferPass(program);
  133. Optimization::TexturePass(env, program);
  134. Optimization::ConstantPropagationPass(program);
  135. Optimization::DeadCodeEliminationPass(program);
  136. Optimization::IdentityRemovalPass(program);
  137. Optimization::VerificationPass(program);
  138. Optimization::CollectShaderInfoPass(env, program);
  139. CollectInterpolationInfo(env, program);
  140. AddNVNStorageBuffers(program);
  141. return program;
  142. }
  143. } // namespace Shader::Maxwell