AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Add an option to SPIRV-Cross to enforce invariant floating point math to prevent different depth calculation between prepass & basepass when running on Metal 2.0 and earlier.

Browse Source
This commit is contained in:
Mark Satterthwaite 2019-08-14 11:25:18 -04:00 committed by Lukas Hermanns
parent e4c6388571
commit 69b703f1da
2 changed files with 60 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

62
spirv_msl.cpp
View File

@ -977,7 +977,30 @@ void CompilerMSL::preprocess_op_codes()
/* UE Change Begin: Allow Metal to use the array<T> template to make arrays a value type */
// UnsafeArray
{
add_header_line(" ");
if (msl_options.invariant_float_math)
{
add_header_line("template <typename T> T fmul(T l, T r) { return metal::fma(l, r, T(0)); }");
add_header_line("template <typename T> T fadd(T l, T r) { return metal::fma(T(1), l, r); }");
add_header_line("template <typename T, int Cols, int Rows> metal::vec<T, Rows> fmul_mv(metal::matrix<T, Cols, Rows> m, metal::vec<T, Cols> v) { metal::vec<T, Rows> res = metal::vec<T, Rows>(0); for(uint i = Cols; i > 0; --i) { res = metal::fma(m[i-1], metal::vec<T, Rows>(v[i-1]), res); } return res; }");
add_header_line("template <typename T, int LCols, int LRows, int RCols, int RRows>");
add_header_line("metal::matrix<T, RCols, LRows> fmul_mat(metal::matrix<T, LCols, LRows> l, metal::matrix<T, RCols, RRows> r)");
add_header_line("{");
add_header_line(" metal::matrix<T, RCols, LRows> res;");
add_header_line(" for(uint i = 0; i < RCols; i++)");
add_header_line(" {");
add_header_line(" metal::vec<T, RCols> tmp(0);");
add_header_line(" for(uint j = 0; j < LCols; j++)");
add_header_line(" {");
add_header_line(" tmp = metal::fma(metal::vec<T, RCols>(r[i][j]),l[j],tmp);");
add_header_line(" }");
add_header_line(" res[i] = tmp;");
add_header_line(" }");
add_header_line(" return res;");
add_header_line("}");
}
add_header_line("template <typename T, size_t Num>");
add_header_line("struct unsafe_array");
add_header_line("{");
@ -5359,6 +5382,20 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
case OpFRem:
MSL_BFOP(fmod);
break;
case OpFMul:
if (msl_options.invariant_float_math)
MSL_BFOP(fmul);
else
MSL_BOP(*);
break;
case OpFAdd:
if (msl_options.invariant_float_math)
MSL_BFOP(fadd);
else
MSL_BOP(+);
break;
// Atomics
case OpAtomicExchange:
@ -5781,15 +5818,32 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
uint32_t id = ops[1];
uint32_t a = ops[2];
uint32_t b = ops[3];
auto *type_a = maybe_get<SPIRType>(a);
auto *type_b = maybe_get<SPIRType>(b);
auto &type = get<SPIRType>(result_type);
string expr = type_to_glsl_constructor(type);
expr += "(";
for (uint32_t col = 0; col < type.columns; col++)
{
expr += to_enclosed_expression(a);
expr += " * ";
expr += to_extract_component_expression(b, col);
if (msl_options.invariant_float_math && type_a && type_b &&
( (is_matrix(*type_a) && is_matrix(*type_b)) ||
(is_matrix(*type_a) && is_vector(*type_b)) ||
(is_vector(*type_a) && is_matrix(*type_b)) ) )
{
expr += "fmul_mat(";
expr += to_enclosed_expression(a);
expr += ", ";
expr += to_extract_component_expression(b, col);
expr += ")";
}
else
{
expr += to_enclosed_expression(a);
expr += " * ";
expr += to_extract_component_expression(b, col);
}
if (col + 1 < type.columns)
expr += ", ";
}

2
spirv_msl.hpp
View File

@ -299,6 +299,8 @@ public:
bool enforce_storge_buffer_bounds = false;
/* UE Change End: Storage buffer robustness - clamps access to SSBOs to the size of the buffer */
bool invariant_float_math = false;
// Requires MSL 2.1, use the native support for texel buffers.
bool texture_buffer_native = false;