emit_context.cpp 4.3 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113
  1. // Copyright 2021 yuzu Emulator Project
  2. // Licensed under GPLv2 or any later version
  3. // Refer to the license.txt file included.
  4. #include "shader_recompiler/backend/bindings.h"
  5. #include "shader_recompiler/backend/glsl/emit_context.h"
  6. #include "shader_recompiler/frontend/ir/program.h"
  7. namespace Shader::Backend::GLSL {
  8. EmitContext::EmitContext(IR::Program& program, [[maybe_unused]] Bindings& bindings,
  9. const Profile& profile_)
  10. : info{program.info}, profile{profile_} {
  11. std::string header = "#version 450\n";
  12. SetupExtensions(header);
  13. if (program.stage == Stage::Compute) {
  14. header += fmt::format("layout(local_size_x={},local_size_y={},local_size_z={}) in;\n",
  15. program.workgroup_size[0], program.workgroup_size[1],
  16. program.workgroup_size[2]);
  17. }
  18. code += header;
  19. DefineConstantBuffers();
  20. DefineStorageBuffers();
  21. DefineHelperFunctions();
  22. code += "void main(){\n";
  23. }
  24. void EmitContext::SetupExtensions(std::string& header) {
  25. if (info.uses_int64) {
  26. header += "#extension GL_ARB_gpu_shader_int64 : enable\n";
  27. }
  28. if (info.uses_int64_bit_atomics) {
  29. header += "#extension GL_NV_shader_atomic_int64 : enable\n";
  30. }
  31. if (info.uses_atomic_f32_add) {
  32. header += "#extension GL_NV_shader_atomic_float : enable\n";
  33. }
  34. if (info.uses_atomic_f16x2_add || info.uses_atomic_f16x2_min || info.uses_atomic_f16x2_max) {
  35. header += "#extension NV_shader_atomic_fp16_vector : enable\n";
  36. }
  37. if (info.uses_fp16) {
  38. // TODO: AMD
  39. header += "#extension GL_NV_gpu_shader5 : enable\n";
  40. }
  41. }
  42. void EmitContext::DefineConstantBuffers() {
  43. if (info.constant_buffer_descriptors.empty()) {
  44. return;
  45. }
  46. u32 binding{};
  47. for (const auto& desc : info.constant_buffer_descriptors) {
  48. Add("layout(std140,binding={}) uniform cbuf_{}{{vec4 cbuf{}[{}];}};", binding, binding,
  49. desc.index, 4 * 1024);
  50. ++binding;
  51. }
  52. }
  53. void EmitContext::DefineStorageBuffers() {
  54. if (info.storage_buffers_descriptors.empty()) {
  55. return;
  56. }
  57. u32 binding{};
  58. for (const auto& desc : info.storage_buffers_descriptors) {
  59. Add("layout(std430,binding={}) buffer ssbo_{}{{uint ssbo{}[];}};", binding, binding,
  60. desc.cbuf_index, desc.count);
  61. ++binding;
  62. }
  63. }
  64. void EmitContext::DefineHelperFunctions() {
  65. if (info.uses_global_increment) {
  66. code += "uint CasIncrement(uint op_a,uint op_b){return(op_a>=op_b)?0u:(op_a+1u);}\n";
  67. }
  68. if (info.uses_global_decrement) {
  69. code +=
  70. "uint CasDecrement(uint op_a,uint op_b){return(op_a==0||op_a>op_b)?op_b:(op_a-1u);}\n";
  71. }
  72. if (info.uses_atomic_f32_add) {
  73. code += "uint CasFloatAdd(uint op_a,uint op_b){return "
  74. "floatBitsToUint(uintBitsToFloat(op_a)+uintBitsToFloat(op_b));}\n";
  75. }
  76. if (info.uses_atomic_f32x2_add) {
  77. code += "uint CasFloatAdd32x2(uint op_a,uint op_b){return "
  78. "packHalf2x16(unpackHalf2x16(op_a)+unpackHalf2x16(op_b));}\n";
  79. }
  80. if (info.uses_atomic_f32x2_min) {
  81. code += "uint CasFloatMin32x2(uint op_a,uint op_b){return "
  82. "packHalf2x16(min(unpackHalf2x16(op_a),unpackHalf2x16(op_b)));}\n";
  83. }
  84. if (info.uses_atomic_f32x2_max) {
  85. code += "uint CasFloatMax32x2(uint op_a,uint op_b){return "
  86. "packHalf2x16(max(unpackHalf2x16(op_a),unpackHalf2x16(op_b)));}\n";
  87. }
  88. if (info.uses_atomic_f16x2_add) {
  89. code += "uint CasFloatAdd16x2(uint op_a,uint op_b){return "
  90. "packFloat2x16(unpackFloat2x16(op_a)+unpackFloat2x16(op_b));}\n";
  91. }
  92. if (info.uses_atomic_f16x2_min) {
  93. code += "uint CasFloatMin16x2(uint op_a,uint op_b){return "
  94. "packFloat2x16(min(unpackFloat2x16(op_a),unpackFloat2x16(op_b)));}\n";
  95. }
  96. if (info.uses_atomic_f16x2_max) {
  97. code += "uint CasFloatMax16x2(uint op_a,uint op_b){return "
  98. "packFloat2x16(max(unpackFloat2x16(op_a),unpackFloat2x16(op_b)));}\n";
  99. }
  100. // TODO: Track this usage
  101. code += "uint CasMinS32(uint op_a,uint op_b){return uint(min(int(op_a),int(op_b)));}";
  102. code += "uint CasMaxS32(uint op_a,uint op_b){return uint(max(int(op_a),int(op_b)));}";
  103. }
  104. } // namespace Shader::Backend::GLSL