mirror of
https://git.suyu.dev/suyu/suyu.git
synced 2026-02-10 09:12:00 +00:00
397 lines
16 KiB
C++
397 lines
16 KiB
C++
// SPDX-FileCopyrightText: Copyright 2024 suyu Emulator Project
|
|
// SPDX-License-Identifier: GPL-2.0-or-later
|
|
|
|
#include "common/div_ceil.h"
|
|
#include "shader_recompiler/backend/bindings.h"
|
|
#include "shader_recompiler/backend/msl/msl_emit_context.h"
|
|
#include "shader_recompiler/frontend/ir/program.h"
|
|
#include "shader_recompiler/profile.h"
|
|
#include "shader_recompiler/runtime_info.h"
|
|
|
|
namespace Shader::Backend::MSL {
|
|
namespace {
|
|
u32 CbufIndex(size_t offset) {
|
|
return (offset / 4) % 4;
|
|
}
|
|
|
|
char Swizzle(size_t offset) {
|
|
return "xyzw"[CbufIndex(offset)];
|
|
}
|
|
|
|
std::string_view InterpDecorator(Interpolation interp) {
|
|
switch (interp) {
|
|
case Interpolation::Smooth:
|
|
return "";
|
|
case Interpolation::Flat:
|
|
return "flat ";
|
|
case Interpolation::NoPerspective:
|
|
return "noperspective ";
|
|
}
|
|
throw InvalidArgument("Invalid interpolation {}", interp);
|
|
}
|
|
|
|
// TODO
|
|
std::string_view InputArrayDecorator(Stage stage) {
|
|
switch (stage) {
|
|
case Stage::Geometry:
|
|
case Stage::TessellationControl:
|
|
case Stage::TessellationEval:
|
|
return "[]";
|
|
default:
|
|
return "";
|
|
}
|
|
}
|
|
|
|
// TODO
|
|
std::string OutputDecorator(Stage stage, u32 size) {
|
|
switch (stage) {
|
|
case Stage::TessellationControl:
|
|
return fmt::format("[{}]", size);
|
|
default:
|
|
return "";
|
|
}
|
|
}
|
|
|
|
// TODO
|
|
std::string_view GetTessMode(TessPrimitive primitive) {
|
|
switch (primitive) {
|
|
case TessPrimitive::Triangles:
|
|
return "triangles";
|
|
case TessPrimitive::Quads:
|
|
return "quads";
|
|
case TessPrimitive::Isolines:
|
|
return "isolines";
|
|
}
|
|
throw InvalidArgument("Invalid tessellation primitive {}", primitive);
|
|
}
|
|
|
|
// TODO
|
|
std::string_view GetTessSpacing(TessSpacing spacing) {
|
|
switch (spacing) {
|
|
case TessSpacing::Equal:
|
|
return "equal_spacing";
|
|
case TessSpacing::FractionalOdd:
|
|
return "fractional_odd_spacing";
|
|
case TessSpacing::FractionalEven:
|
|
return "fractional_even_spacing";
|
|
}
|
|
throw InvalidArgument("Invalid tessellation spacing {}", spacing);
|
|
}
|
|
|
|
// TODO
|
|
std::string_view InputPrimitive(InputTopology topology) {
|
|
switch (topology) {
|
|
case InputTopology::Points:
|
|
return "points";
|
|
case InputTopology::Lines:
|
|
return "lines";
|
|
case InputTopology::LinesAdjacency:
|
|
return "lines_adjacency";
|
|
case InputTopology::Triangles:
|
|
return "triangles";
|
|
case InputTopology::TrianglesAdjacency:
|
|
return "triangles_adjacency";
|
|
}
|
|
throw InvalidArgument("Invalid input topology {}", topology);
|
|
}
|
|
|
|
// TODO
|
|
std::string_view OutputPrimitive(OutputTopology topology) {
|
|
switch (topology) {
|
|
case OutputTopology::PointList:
|
|
return "points";
|
|
case OutputTopology::LineStrip:
|
|
return "line_strip";
|
|
case OutputTopology::TriangleStrip:
|
|
return "triangle_strip";
|
|
}
|
|
throw InvalidArgument("Invalid output topology {}", topology);
|
|
}
|
|
} // Anonymous namespace
|
|
|
|
EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_,
|
|
const RuntimeInfo& runtime_info_)
|
|
: info{program.info}, profile{profile_}, runtime_info{runtime_info_}, stage{program.stage},
|
|
uses_geometry_passthrough{program.is_geometry_passthrough &&
|
|
profile.support_geometry_shader_passthrough} {
|
|
if (profile.need_fastmath_off) {
|
|
// TODO
|
|
}
|
|
switch (program.stage) {
|
|
case Stage::VertexA:
|
|
case Stage::VertexB:
|
|
stage_name = "vs";
|
|
break;
|
|
case Stage::TessellationControl:
|
|
stage_name = "tcs";
|
|
header += fmt::format("layout(vertices={})out;", program.invocations);
|
|
break;
|
|
case Stage::TessellationEval:
|
|
stage_name = "tes";
|
|
header += fmt::format("layout({},{},{})in;", GetTessMode(runtime_info.tess_primitive),
|
|
GetTessSpacing(runtime_info.tess_spacing),
|
|
runtime_info.tess_clockwise ? "cw" : "ccw");
|
|
break;
|
|
case Stage::Geometry:
|
|
stage_name = "gs";
|
|
header += fmt::format("layout({})in;", InputPrimitive(runtime_info.input_topology));
|
|
if (uses_geometry_passthrough) {
|
|
header += "layout(passthrough)in gl_PerVertex{vec4 gl_Position;};";
|
|
break;
|
|
} else if (program.is_geometry_passthrough &&
|
|
!profile.support_geometry_shader_passthrough) {
|
|
LOG_WARNING(Shader_MSL, "Passthrough geometry program used but not supported");
|
|
}
|
|
header += fmt::format(
|
|
"layout({},max_vertices={})out;in gl_PerVertex{{vec4 gl_Position;}}gl_in[];",
|
|
OutputPrimitive(program.output_topology), program.output_vertices);
|
|
break;
|
|
case Stage::Fragment:
|
|
stage_name = "fs";
|
|
if (runtime_info.force_early_z) {
|
|
header += "layout(early_fragment_tests)in;";
|
|
}
|
|
break;
|
|
case Stage::Compute:
|
|
stage_name = "cs";
|
|
const u32 local_x{std::max(program.workgroup_size[0], 1u)};
|
|
const u32 local_y{std::max(program.workgroup_size[1], 1u)};
|
|
const u32 local_z{std::max(program.workgroup_size[2], 1u)};
|
|
header += fmt::format("layout(local_size_x={},local_size_y={},local_size_z={}) in;",
|
|
local_x, local_y, local_z);
|
|
break;
|
|
}
|
|
// TODO
|
|
// SetupOutPerVertex(*this, header);
|
|
// SetupInPerVertex(*this, header);
|
|
|
|
for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
|
|
if (!info.loads.Generic(index) || !runtime_info.previous_stage_stores.Generic(index)) {
|
|
continue;
|
|
}
|
|
const auto qualifier{uses_geometry_passthrough ? "passthrough"
|
|
: fmt::format("location={}", index)};
|
|
header += fmt::format("layout({}){}in vec4 in_attr{}{};", qualifier,
|
|
InterpDecorator(info.interpolation[index]), index,
|
|
InputArrayDecorator(stage));
|
|
}
|
|
for (size_t index = 0; index < info.uses_patches.size(); ++index) {
|
|
if (!info.uses_patches[index]) {
|
|
continue;
|
|
}
|
|
const auto qualifier{stage == Stage::TessellationControl ? "out" : "in"};
|
|
header += fmt::format("layout(location={})patch {} vec4 patch{};", index, qualifier, index);
|
|
}
|
|
if (stage == Stage::Fragment) {
|
|
for (size_t index = 0; index < info.stores_frag_color.size(); ++index) {
|
|
if (!info.stores_frag_color[index] && !profile.need_declared_frag_colors) {
|
|
continue;
|
|
}
|
|
header += fmt::format("layout(location={})out vec4 frag_color{};", index, index);
|
|
}
|
|
}
|
|
for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
|
|
if (info.stores.Generic(index)) {
|
|
DefineGenericOutput(index, program.invocations);
|
|
}
|
|
}
|
|
if (info.uses_rescaling_uniform) {
|
|
header += "layout(location=0) uniform vec4 scaling;";
|
|
}
|
|
if (info.uses_render_area) {
|
|
header += "layout(location=1) uniform vec4 render_area;";
|
|
}
|
|
DefineHelperFunctions();
|
|
DefineConstants();
|
|
}
|
|
|
|
// TODO
|
|
void EmitContext::DefineGenericOutput(size_t index, u32 invocations) {
|
|
static constexpr std::string_view swizzle{"xyzw"};
|
|
const size_t base_index{static_cast<size_t>(IR::Attribute::Generic0X) + index * 4};
|
|
u32 element{0};
|
|
while (element < 4) {
|
|
std::string definition{fmt::format("layout(location={}", index)};
|
|
const u32 remainder{4 - element};
|
|
const TransformFeedbackVarying* xfb_varying{};
|
|
const size_t xfb_varying_index{base_index + element};
|
|
if (xfb_varying_index < runtime_info.xfb_count) {
|
|
xfb_varying = &runtime_info.xfb_varyings[xfb_varying_index];
|
|
xfb_varying = xfb_varying->components > 0 ? xfb_varying : nullptr;
|
|
}
|
|
const u32 num_components{xfb_varying ? xfb_varying->components : remainder};
|
|
if (element > 0) {
|
|
definition += fmt::format(",component={}", element);
|
|
}
|
|
if (xfb_varying) {
|
|
definition +=
|
|
fmt::format(",xfb_buffer={},xfb_stride={},xfb_offset={}", xfb_varying->buffer,
|
|
xfb_varying->stride, xfb_varying->offset);
|
|
}
|
|
std::string name{fmt::format("out_attr{}", index)};
|
|
if (num_components < 4 || element > 0) {
|
|
name += fmt::format("_{}", swizzle.substr(element, num_components));
|
|
}
|
|
const auto type{num_components == 1 ? "float" : fmt::format("vec{}", num_components)};
|
|
definition += fmt::format(")out {} {}{};", type, name, OutputDecorator(stage, invocations));
|
|
header += definition;
|
|
|
|
const GenericElementInfo element_info{
|
|
.name = name,
|
|
.first_element = element,
|
|
.num_components = num_components,
|
|
};
|
|
std::fill_n(output_generics[index].begin() + element, num_components, element_info);
|
|
element += num_components;
|
|
}
|
|
}
|
|
|
|
void EmitContext::DefineHelperFunctions() {
|
|
header += "\n#define ftoi floatBitsToInt\n#define ftou floatBitsToUint\n"
|
|
"#define itof intBitsToFloat\n#define utof uintBitsToFloat\n";
|
|
if (info.uses_global_increment || info.uses_shared_increment) {
|
|
header += "uint CasIncrement(uint op_a,uint op_b){return op_a>=op_b?0u:(op_a+1u);}";
|
|
}
|
|
if (info.uses_global_decrement || info.uses_shared_decrement) {
|
|
header += "uint CasDecrement(uint op_a,uint op_b){"
|
|
"return op_a==0||op_a>op_b?op_b:(op_a-1u);}";
|
|
}
|
|
if (info.uses_atomic_f32_add) {
|
|
header += "uint CasFloatAdd(uint op_a,float op_b){"
|
|
"return ftou(utof(op_a)+op_b);}";
|
|
}
|
|
if (info.uses_atomic_f32x2_add) {
|
|
header += "uint CasFloatAdd32x2(uint op_a,vec2 op_b){"
|
|
"return packHalf2x16(unpackHalf2x16(op_a)+op_b);}";
|
|
}
|
|
if (info.uses_atomic_f32x2_min) {
|
|
header += "uint CasFloatMin32x2(uint op_a,vec2 op_b){return "
|
|
"packHalf2x16(min(unpackHalf2x16(op_a),op_b));}";
|
|
}
|
|
if (info.uses_atomic_f32x2_max) {
|
|
header += "uint CasFloatMax32x2(uint op_a,vec2 op_b){return "
|
|
"packHalf2x16(max(unpackHalf2x16(op_a),op_b));}";
|
|
}
|
|
if (info.uses_atomic_f16x2_add) {
|
|
header += "uint CasFloatAdd16x2(uint op_a,f16vec2 op_b){return "
|
|
"packFloat2x16(unpackFloat2x16(op_a)+op_b);}";
|
|
}
|
|
if (info.uses_atomic_f16x2_min) {
|
|
header += "uint CasFloatMin16x2(uint op_a,f16vec2 op_b){return "
|
|
"packFloat2x16(min(unpackFloat2x16(op_a),op_b));}";
|
|
}
|
|
if (info.uses_atomic_f16x2_max) {
|
|
header += "uint CasFloatMax16x2(uint op_a,f16vec2 op_b){return "
|
|
"packFloat2x16(max(unpackFloat2x16(op_a),op_b));}";
|
|
}
|
|
if (info.uses_atomic_s32_min) {
|
|
header += "uint CasMinS32(uint op_a,uint op_b){return uint(min(int(op_a),int(op_b)));}";
|
|
}
|
|
if (info.uses_atomic_s32_max) {
|
|
header += "uint CasMaxS32(uint op_a,uint op_b){return uint(max(int(op_a),int(op_b)));}";
|
|
}
|
|
if (info.uses_global_memory && profile.support_int64) {
|
|
header += DefineGlobalMemoryFunctions();
|
|
}
|
|
if (info.loads_indexed_attributes) {
|
|
const bool is_array{stage == Stage::Geometry};
|
|
const auto vertex_arg{is_array ? ",uint vertex" : ""};
|
|
std::string func{
|
|
fmt::format("float IndexedAttrLoad(int offset{}){{int base_index=offset>>2;uint "
|
|
"masked_index=uint(base_index)&3u;switch(base_index>>2){{",
|
|
vertex_arg)};
|
|
if (info.loads.AnyComponent(IR::Attribute::PositionX)) {
|
|
const auto position_idx{is_array ? "gl_in[vertex]." : ""};
|
|
func += fmt::format("case {}:return {}{}[masked_index];",
|
|
static_cast<u32>(IR::Attribute::PositionX) >> 2, position_idx,
|
|
"__out.position");
|
|
}
|
|
const u32 base_attribute_value = static_cast<u32>(IR::Attribute::Generic0X) >> 2;
|
|
for (u32 index = 0; index < IR::NUM_GENERICS; ++index) {
|
|
if (!info.loads.Generic(index)) {
|
|
continue;
|
|
}
|
|
const auto vertex_idx{is_array ? "[vertex]" : ""};
|
|
func += fmt::format("case {}:return in_attr{}{}[masked_index];",
|
|
base_attribute_value + index, index, vertex_idx);
|
|
}
|
|
func += "default: return 0.0;}}";
|
|
header += func;
|
|
}
|
|
if (info.stores_indexed_attributes) {
|
|
// TODO
|
|
}
|
|
}
|
|
|
|
std::string EmitContext::DefineGlobalMemoryFunctions() {
|
|
const auto define_body{[&](std::string& func, size_t index, std::string_view return_statement) {
|
|
const auto& ssbo{info.storage_buffers_descriptors[index]};
|
|
const u32 size_cbuf_offset{ssbo.cbuf_offset + 8};
|
|
const auto ssbo_addr{fmt::format("ssbo_addr{}", index)};
|
|
const auto cbuf{fmt::format("{}_cbuf{}", stage_name, ssbo.cbuf_index)};
|
|
std::array<std::string, 2> addr_xy;
|
|
std::array<std::string, 2> size_xy;
|
|
for (size_t i = 0; i < addr_xy.size(); ++i) {
|
|
const auto addr_loc{ssbo.cbuf_offset + 4 * i};
|
|
const auto size_loc{size_cbuf_offset + 4 * i};
|
|
addr_xy[i] = fmt::format("ftou({}[{}].{})", cbuf, addr_loc / 16, Swizzle(addr_loc));
|
|
size_xy[i] = fmt::format("ftou({}[{}].{})", cbuf, size_loc / 16, Swizzle(size_loc));
|
|
}
|
|
const u32 ssbo_align_mask{~(static_cast<u32>(profile.min_ssbo_alignment) - 1U)};
|
|
const auto aligned_low_addr{fmt::format("{}&{}", addr_xy[0], ssbo_align_mask)};
|
|
const auto aligned_addr{fmt::format("uvec2({},{})", aligned_low_addr, addr_xy[1])};
|
|
const auto addr_pack{fmt::format("packUint2x32({})", aligned_addr)};
|
|
const auto addr_statement{fmt::format("uint64_t {}={};", ssbo_addr, addr_pack)};
|
|
func += addr_statement;
|
|
|
|
const auto size_vec{fmt::format("uvec2({},{})", size_xy[0], size_xy[1])};
|
|
const auto comp_lhs{fmt::format("(addr>={})", ssbo_addr)};
|
|
const auto comp_rhs{fmt::format("(addr<({}+uint64_t({})))", ssbo_addr, size_vec)};
|
|
const auto comparison{fmt::format("if({}&&{}){{", comp_lhs, comp_rhs)};
|
|
func += comparison;
|
|
|
|
const auto ssbo_name{fmt::format("{}_ssbo{}", stage_name, index)};
|
|
func += fmt::format(fmt::runtime(return_statement), ssbo_name, ssbo_addr);
|
|
}};
|
|
std::string write_func{"void WriteGlobal32(uint64_t addr,uint data){"};
|
|
std::string write_func_64{"void WriteGlobal64(uint64_t addr,uvec2 data){"};
|
|
std::string write_func_128{"void WriteGlobal128(uint64_t addr,uvec4 data){"};
|
|
std::string load_func{"uint LoadGlobal32(uint64_t addr){"};
|
|
std::string load_func_64{"uvec2 LoadGlobal64(uint64_t addr){"};
|
|
std::string load_func_128{"uvec4 LoadGlobal128(uint64_t addr){"};
|
|
const size_t num_buffers{info.storage_buffers_descriptors.size()};
|
|
for (size_t index = 0; index < num_buffers; ++index) {
|
|
if (!info.nvn_buffer_used[index]) {
|
|
continue;
|
|
}
|
|
define_body(write_func, index, "{0}[uint(addr-{1})>>2]=data;return;}}");
|
|
define_body(write_func_64, index,
|
|
"{0}[uint(addr-{1})>>2]=data.x;{0}[uint(addr-{1}+4)>>2]=data.y;return;}}");
|
|
define_body(write_func_128, index,
|
|
"{0}[uint(addr-{1})>>2]=data.x;{0}[uint(addr-{1}+4)>>2]=data.y;{0}[uint("
|
|
"addr-{1}+8)>>2]=data.z;{0}[uint(addr-{1}+12)>>2]=data.w;return;}}");
|
|
define_body(load_func, index, "return {0}[uint(addr-{1})>>2];}}");
|
|
define_body(load_func_64, index,
|
|
"return uvec2({0}[uint(addr-{1})>>2],{0}[uint(addr-{1}+4)>>2]);}}");
|
|
define_body(load_func_128, index,
|
|
"return uvec4({0}[uint(addr-{1})>>2],{0}[uint(addr-{1}+4)>>2],{0}["
|
|
"uint(addr-{1}+8)>>2],{0}[uint(addr-{1}+12)>>2]);}}");
|
|
}
|
|
write_func += '}';
|
|
write_func_64 += '}';
|
|
write_func_128 += '}';
|
|
load_func += "return 0u;}";
|
|
load_func_64 += "return uint2(0);}";
|
|
load_func_128 += "return uint4(0);}";
|
|
return write_func + write_func_64 + write_func_128 + load_func + load_func_64 + load_func_128;
|
|
}
|
|
|
|
void EmitContext::DefineConstants() {
|
|
if (info.uses_fswzadd) {
|
|
header += "const float FSWZ_A[]=float[4](-1.f,1.f,-1.f,0.f);"
|
|
"const float FSWZ_B[]=float[4](-1.f,-1.f,1.f,-1.f);";
|
|
}
|
|
}
|
|
|
|
} // namespace Shader::Backend::MSL
|