shadps4-emu/shadPS4
1
0
Fork 0
mirror of https://github.com/shadps4-emu/shadPS4.git synced 2026-10-12 02:22:27 +00:00
Code Issues Packages Projects Releases Wiki Activity

shader_recompiler: Improve NaN handling (#5257)

Browse Source 
* translate: Handle NaN semantics of legacy instruction variants

* translate: Use unordered float compares for v_cmp_n* instructions

The ACO compiler in mesa emits those variants for unordered compares so use those to get correct NaN behavior

* spirv_emit_float_point: Proper NaN semantics for three number compare instructions

Manual mentions DX10 NaN handling for all of them and for V_MED3_F32 specifically if any input is NaN it fallbacks to MIN3 instruction so implement those properly using NMin/NMax

* translate: Honor DX10_CLAMP mode bit

Switch FPSaturate to using NClamp instead of FClamp, to match HLSL saturate which returns 0.f with NaN input. Otherwise use normal FClamp which propagates the NaN instead

* emit_spirv_floating_pont: Dont use amd extension
This commit is contained in:
TheTurtle
2026-10-04 13:37:56 -07:00
committed by GitHub
parent 16e79a4296
commit c495290879
11 changed files with 140 additions and 102 deletions
Download Patch File Download Diff File Expand all files Collapse all files
externals/sirit
Vendored
+1 -1
Submodule externals/sirit updated: c58f4d441c...fc2b42d5e5
src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp
+10 -16
View File
@@ -56,25 +56,21 @@ Id EmitFPMin64(EmitContext& ctx, Id a, Id b) {
}
Id EmitFPMinTri32(EmitContext& ctx, Id a, Id b, Id c) {
if (ctx.profile.supports_trinary_minmax) {
return ctx.OpFMin3AMD(ctx.F32[1], a, b, c);
}
return ctx.OpFMin(ctx.F32[1], a, ctx.OpFMin(ctx.F32[1], b, c));
return ctx.OpNMin(ctx.F32[1], a, ctx.OpNMin(ctx.F32[1], b, c));
}
Id EmitFPMaxTri32(EmitContext& ctx, Id a, Id b, Id c) {
if (ctx.profile.supports_trinary_minmax) {
return ctx.OpFMax3AMD(ctx.F32[1], a, b, c);
}
return ctx.OpFMax(ctx.F32[1], a, ctx.OpFMax(ctx.F32[1], b, c));
return ctx.OpNMax(ctx.F32[1], a, ctx.OpNMax(ctx.F32[1], b, c));
}
Id EmitFPMedTri32(EmitContext& ctx, Id a, Id b, Id c) {
if (ctx.profile.supports_trinary_minmax) {
return ctx.OpFMid3AMD(ctx.F32[1], a, b, c);
}
const Id is_nan{ctx.OpLogicalOr(
ctx.U1[1], ctx.OpIsNan(ctx.U1[1], a),
ctx.OpLogicalOr(ctx.U1[1], ctx.OpIsNan(ctx.U1[1], b), ctx.OpIsNan(ctx.U1[1], c)))};
const Id mmx{ctx.OpFMin(ctx.F32[1], ctx.OpFMax(ctx.F32[1], a, b), c)};
return ctx.OpFMax(ctx.F32[1], ctx.OpFMin(ctx.F32[1], a, b), mmx);
const Id med3{ctx.OpFMax(ctx.F32[1], ctx.OpFMin(ctx.F32[1], a, b), mmx)};
const Id min3{ctx.OpNMin(ctx.F32[1], a, ctx.OpNMin(ctx.F32[1], b, c))};
return ctx.OpSelect(ctx.F32[1], is_nan, min3, med3);
}
Id EmitFPMul32(EmitContext& ctx, IR::Inst* inst, Id a, Id b) {
@@ -146,15 +142,13 @@ Id EmitFPSqrt(EmitContext& ctx, Id value) {
}
Id EmitFPSaturate32(EmitContext& ctx, Id value) {
const Id zero{ctx.ConstF32(f32{0.0})};
const Id one{ctx.ConstF32(f32{1.0})};
return ctx.OpFClamp(ctx.F32[1], value, zero, one);
return ctx.OpNClamp(ctx.F32[1], value, ctx.f32_zero_value, ctx.f32_one_value);
}
Id EmitFPSaturate64(EmitContext& ctx, Id value) {
const Id zero{ctx.Constant(ctx.F64[1], f64{0.0})};
const Id one{ctx.Constant(ctx.F64[1], f64{1.0})};
return ctx.OpFClamp(ctx.F64[1], value, zero, one);
return ctx.OpNClamp(ctx.F64[1], value, zero, one);
}
Id EmitFPClamp32(EmitContext& ctx, Id value, Id min_value, Id max_value) {
src/shader_recompiler/backend/spirv/spirv_emit_context.cpp
+2 -1
View File
@@ -152,7 +152,8 @@ void EmitContext::DefineArithmeticTypes() {
u16_zero_value = Constant(U16, 0U);
u32_one_value = ConstU32(1U);
u32_zero_value = ConstU32(0U);
f32_zero_value = ConstF32(0.0f);
f32_one_value = ConstF32(1.f);
f32_zero_value = ConstF32(0.f);
u64_one_value = Constant(U64, 1ULL);
u64_zero_value = Constant(U64, 0ULL);
src/shader_recompiler/backend/spirv/spirv_emit_context.h
+1
View File
@@ -225,6 +225,7 @@ public:
Id u16_zero_value{};
Id u32_one_value{};
Id u32_zero_value{};
Id f32_one_value{};
Id f32_zero_value{};
Id u64_one_value{};
Id u64_zero_value{};
src/shader_recompiler/frontend/translate/translate.cpp
+21 -4
View File
@@ -973,7 +973,11 @@ void Translator::SetDst(const InstOperand& operand, const IR::U32F32& value) {
result = ir.FPMul(result, ir.Imm32(operand.output_modifier.multiplier));
}
if (operand.output_modifier.clamp) {
result = ir.FPSaturate(result);
if (runtime_info.props.dx10_clamp) {
result = ir.FPSaturate(result);
} else {
result = ir.FPClamp(result, ir.Imm32(0.f), ir.Imm32(1.f));
}
}
}
@@ -1001,7 +1005,11 @@ void Translator::SetDst16(const InstOperand& operand, const IR::U32F32& value) {
result = ir.FPMul(result, ir.Imm32(operand.output_modifier.multiplier));
}
if (operand.output_modifier.clamp) {
result = ir.FPSaturate(result);
if (runtime_info.props.dx10_clamp) {
result = ir.FPSaturate(result);
} else {
result = ir.FPClamp(result, ir.Imm32(0.f), ir.Imm32(1.f));
}
}
} else {
if (operand.output_modifier.clamp) {
@@ -1058,7 +1066,11 @@ void Translator::SetDst64(const InstOperand& operand, const IR::U64F64& value_ra
ir.FPMul(value_untyped, ir.Imm64(f64(operand.output_modifier.multiplier)));
}
if (operand.output_modifier.clamp) {
value_untyped = ir.FPSaturate(value_untyped);
if (runtime_info.props.dx10_clamp) {
value_untyped = ir.FPSaturate(value_untyped);
} else {
value_untyped = ir.FPClamp(value_untyped, ir.Imm32(0.f), ir.Imm32(1.f));
}
}
}
@@ -1102,7 +1114,12 @@ void Translator::SetDstPk(const InstOperand& operand, const pk_type<T>& value) {
if constexpr (std::is_same_v<T, IR::F32>) {
if (operand.output_modifier.clamp) {
v = {ir.FPSaturate(v.first), ir.FPSaturate(v.second)};
if (runtime_info.props.dx10_clamp) {
v = {ir.FPSaturate(v.first), ir.FPSaturate(v.second)};
} else {
v = {ir.FPClamp(v.first, ir.Imm32(0.f), ir.Imm32(1.f)),
ir.FPClamp(v.second, ir.Imm32(0.f), ir.Imm32(1.f))};
}
}
} else {
if (operand.output_modifier.clamp) {
src/shader_recompiler/frontend/translate/translate.h
+5 -2
View File
@@ -162,6 +162,8 @@ public:
void V_SUBREV_F32(const GcnInst& inst);
void V_MUL_F32(const GcnInst& inst);
void V_MUL_LEGACY_F32(const GcnInst& inst);
void V_MAC_LEGACY_F32(const GcnInst& inst);
void V_MAD_LEGACY_F32(const GcnInst& inst);
void V_MUL_I32_I24(const GcnInst& inst, bool is_signed);
void V_MIN_F32(const GcnInst& inst, bool is_legacy = false);
void V_MAX_F32(const GcnInst& inst, bool is_legacy = false);
@@ -246,8 +248,8 @@ public:
void V_MOVRELSD_B32(const GcnInst& inst);
// VOPC
void V_CMP_F32(ConditionOp op, bool set_exec, const GcnInst& inst);
void V_CMP_F64(ConditionOp op, bool set_exec, const GcnInst& inst);
void V_CMP_F32(ConditionOp op, bool ordered, bool set_exec, const GcnInst& inst);
void V_CMP_F64(ConditionOp op, bool ordered, bool set_exec, const GcnInst& inst);
void V_CMP_U32(ConditionOp op, bool is_signed, bool set_exec, const GcnInst& inst);
void V_CMP_U64(ConditionOp op, bool is_signed, bool set_exec, const GcnInst& inst);
void V_CMP_CLASS_F32(const GcnInst& inst);
@@ -375,6 +377,7 @@ private:
// Vector ALU Helpers
IR::U32 GetCarryIn(const GcnInst& inst);
void SetCarryOut(const GcnInst& inst, const IR::U1& carry);
IR::F32 LegacyMul(const IR::F32& a, const IR::F32& b);
IR::U32 VMovRelSHelper(u32 src_vgprno, const IR::U32 m0);
void VMovRelDHelper(u32 dst_vgprno, const IR::U32 src_val, const IR::U32 m0);
src/shader_recompiler/frontend/translate/vector_alu.cpp
+77 -65
View File
@@ -24,7 +24,7 @@ void Translator::EmitVectorAlu(const GcnInst& inst) {
case Opcode::V_SUBREV_F32:
return V_SUBREV_F32(inst);
case Opcode::V_MAC_LEGACY_F32:
return V_MAC_F32(inst);
return V_MAC_LEGACY_F32(inst);
case Opcode::V_MUL_LEGACY_F32:
return V_MUL_LEGACY_F32(inst);
case Opcode::V_MUL_F32:
@@ -225,89 +225,87 @@ void Translator::EmitVectorAlu(const GcnInst& inst) {
// VOPC
// V_CMP_{OP16}_F32
case Opcode::V_CMP_F_F32:
return V_CMP_F32(ConditionOp::F, false, inst);
return V_CMP_F32(ConditionOp::F, true, false, inst);
case Opcode::V_CMP_LT_F32:
return V_CMP_F32(ConditionOp::LT, false, inst);
return V_CMP_F32(ConditionOp::LT, true, false, inst);
case Opcode::V_CMP_EQ_F32:
return V_CMP_F32(ConditionOp::EQ, false, inst);
return V_CMP_F32(ConditionOp::EQ, true, false, inst);
case Opcode::V_CMP_LE_F32:
return V_CMP_F32(ConditionOp::LE, false, inst);
return V_CMP_F32(ConditionOp::LE, true, false, inst);
case Opcode::V_CMP_GT_F32:
return V_CMP_F32(ConditionOp::GT, false, inst);
return V_CMP_F32(ConditionOp::GT, true, false, inst);
case Opcode::V_CMP_LG_F32:
return V_CMP_F32(ConditionOp::LG, false, inst);
return V_CMP_F32(ConditionOp::LG, true, false, inst);
case Opcode::V_CMP_GE_F32:
return V_CMP_F32(ConditionOp::GE, false, inst);
return V_CMP_F32(ConditionOp::GE, true, false, inst);
case Opcode::V_CMP_U_F32:
return V_CMP_F32(ConditionOp::U, false, inst);
return V_CMP_F32(ConditionOp::U, true, false, inst);
case Opcode::V_CMP_NGE_F32:
return V_CMP_F32(ConditionOp::LT, false, inst);
return V_CMP_F32(ConditionOp::LT, false, false, inst);
case Opcode::V_CMP_NGT_F32:
return V_CMP_F32(ConditionOp::LE, false, inst);
return V_CMP_F32(ConditionOp::LE, false, false, inst);
case Opcode::V_CMP_NLE_F32:
return V_CMP_F32(ConditionOp::GT, false, inst);
return V_CMP_F32(ConditionOp::GT, false, false, inst);
case Opcode::V_CMP_NEQ_F32:
return V_CMP_F32(ConditionOp::LG, false, inst);
return V_CMP_F32(ConditionOp::LG, false, false, inst);
case Opcode::V_CMP_NLT_F32:
return V_CMP_F32(ConditionOp::GE, false, inst);
return V_CMP_F32(ConditionOp::GE, false, false, inst);
// V_CMPX_{OP16}_F32
case Opcode::V_CMPX_F_F32:
return V_CMP_F32(ConditionOp::F, true, inst);
return V_CMP_F32(ConditionOp::F, true, true, inst);
case Opcode::V_CMPX_LT_F32:
return V_CMP_F32(ConditionOp::LT, true, inst);
return V_CMP_F32(ConditionOp::LT, true, true, inst);
case Opcode::V_CMPX_EQ_F32:
return V_CMP_F32(ConditionOp::EQ, true, inst);
return V_CMP_F32(ConditionOp::EQ, true, true, inst);
case Opcode::V_CMPX_LE_F32:
return V_CMP_F32(ConditionOp::LE, true, inst);
return V_CMP_F32(ConditionOp::LE, true, true, inst);
case Opcode::V_CMPX_GT_F32:
return V_CMP_F32(ConditionOp::GT, true, inst);
return V_CMP_F32(ConditionOp::GT, true, true, inst);
case Opcode::V_CMPX_LG_F32:
return V_CMP_F32(ConditionOp::LG, true, inst);
return V_CMP_F32(ConditionOp::LG, true, true, inst);
case Opcode::V_CMPX_GE_F32:
return V_CMP_F32(ConditionOp::GE, true, inst);
return V_CMP_F32(ConditionOp::GE, true, true, inst);
case Opcode::V_CMPX_NGE_F32:
return V_CMP_F32(ConditionOp::LT, true, inst);
return V_CMP_F32(ConditionOp::LT, false, true, inst);
case Opcode::V_CMPX_NLG_F32:
return V_CMP_F32(ConditionOp::EQ, true, inst);
return V_CMP_F32(ConditionOp::EQ, false, true, inst);
case Opcode::V_CMPX_NGT_F32:
return V_CMP_F32(ConditionOp::LE, true, inst);
return V_CMP_F32(ConditionOp::LE, false, true, inst);
case Opcode::V_CMPX_NLE_F32:
return V_CMP_F32(ConditionOp::GT, true, inst);
return V_CMP_F32(ConditionOp::GT, false, true, inst);
case Opcode::V_CMPX_NEQ_F32:
return V_CMP_F32(ConditionOp::LG, true, inst);
return V_CMP_F32(ConditionOp::LG, false, true, inst);
case Opcode::V_CMPX_NLT_F32:
return V_CMP_F32(ConditionOp::GE, true, inst);
case Opcode::V_CMPX_TRU_F32:
return V_CMP_F32(ConditionOp::TRU, true, inst);
return V_CMP_F32(ConditionOp::GE, false, true, inst);
// V_CMP_{OP16}_F64
case Opcode::V_CMP_F_F64:
return V_CMP_F64(ConditionOp::F, false, inst);
return V_CMP_F64(ConditionOp::F, true, false, inst);
case Opcode::V_CMP_LT_F64:
return V_CMP_F64(ConditionOp::LT, false, inst);
return V_CMP_F64(ConditionOp::LT, true, false, inst);
case Opcode::V_CMP_EQ_F64:
return V_CMP_F64(ConditionOp::EQ, false, inst);
return V_CMP_F64(ConditionOp::EQ, true, false, inst);
case Opcode::V_CMP_LE_F64:
return V_CMP_F64(ConditionOp::LE, false, inst);
return V_CMP_F64(ConditionOp::LE, true, false, inst);
case Opcode::V_CMP_GT_F64:
return V_CMP_F64(ConditionOp::GT, false, inst);
return V_CMP_F64(ConditionOp::GT, true, false, inst);
case Opcode::V_CMP_LG_F64:
return V_CMP_F64(ConditionOp::LG, false, inst);
return V_CMP_F64(ConditionOp::LG, true, false, inst);
case Opcode::V_CMP_GE_F64:
return V_CMP_F64(ConditionOp::GE, false, inst);
return V_CMP_F64(ConditionOp::GE, true, false, inst);
case Opcode::V_CMP_U_F64:
return V_CMP_F64(ConditionOp::U, false, inst);
return V_CMP_F64(ConditionOp::U, true, false, inst);
case Opcode::V_CMP_NGE_F64:
return V_CMP_F64(ConditionOp::LT, false, inst);
return V_CMP_F64(ConditionOp::LT, false, false, inst);
case Opcode::V_CMP_NGT_F64:
return V_CMP_F64(ConditionOp::LE, false, inst);
return V_CMP_F64(ConditionOp::LE, false, false, inst);
case Opcode::V_CMP_NLE_F64:
return V_CMP_F64(ConditionOp::GT, false, inst);
return V_CMP_F64(ConditionOp::GT, false, false, inst);
case Opcode::V_CMP_NEQ_F64:
return V_CMP_F64(ConditionOp::LG, false, inst);
return V_CMP_F64(ConditionOp::LG, false, false, inst);
case Opcode::V_CMP_NLT_F64:
return V_CMP_F64(ConditionOp::GE, false, inst);
return V_CMP_F64(ConditionOp::GE, false, false, inst);
// V_CMP_{OP8}_I32
case Opcode::V_CMP_LT_I32:
@@ -416,7 +414,7 @@ void Translator::EmitVectorAlu(const GcnInst& inst) {
// VOP3a
case Opcode::V_MAD_LEGACY_F32:
return V_MAD_F32(inst);
return V_MAD_LEGACY_F32(inst);
case Opcode::V_MAD_F32:
return V_MAD_F32(inst);
case Opcode::V_MAD_I32_I24:
@@ -608,16 +606,23 @@ void Translator::V_MUL_F32(const GcnInst& inst) {
}
void Translator::V_MUL_LEGACY_F32(const GcnInst& inst) {
// GCN V_MUL_LEGACY_F32: if either source is zero, the result is +0.0
// regardless of the other operand (even if NaN or Inf).
// Standard IEEE multiply would produce NaN for 0 * Inf.
const IR::F32 src0{GetSrc<IR::F32>(inst.src[0])};
const IR::F32 src1{GetSrc<IR::F32>(inst.src[1])};
const IR::F32 zero{ir.Imm32(0.0f)};
const IR::U1 src0_zero{ir.FPEqual(src0, zero)};
const IR::U1 src1_zero{ir.FPEqual(src1, zero)};
const IR::U1 either_zero{ir.LogicalOr(src0_zero, src1_zero)};
SetDst(inst.dst[0], IR::F32{ir.Select(either_zero, zero, ir.FPMul(src0, src1))});
SetDst(inst.dst[0], LegacyMul(src0, src1));
}
void Translator::V_MAC_LEGACY_F32(const GcnInst& inst) {
const IR::F32 src0{GetSrc<IR::F32>(inst.src[0])};
const IR::F32 src1{GetSrc<IR::F32>(inst.src[1])};
const IR::F32 dst0{GetSrc<IR::F32>(inst.dst[0])};
SetDst(inst.dst[0], ir.FPAdd(LegacyMul(src0, src1), dst0));
}
void Translator::V_MAD_LEGACY_F32(const GcnInst& inst) {
const IR::F32 src0{GetSrc<IR::F32>(inst.src[0])};
const IR::F32 src1{GetSrc<IR::F32>(inst.src[1])};
const IR::F32 src2{GetSrc<IR::F32>(inst.src[2])};
SetDst(inst.dst[0], ir.FPAdd(LegacyMul(src0, src1), src2));
}
void Translator::V_MUL_I32_I24(const GcnInst& inst, bool is_signed) {
@@ -1181,7 +1186,7 @@ void Translator::V_MOVRELSD_B32(const GcnInst& inst) {
// VOPC
void Translator::V_CMP_F32(ConditionOp op, bool set_exec, const GcnInst& inst) {
void Translator::V_CMP_F32(ConditionOp op, bool ordered, bool set_exec, const GcnInst& inst) {
const IR::F32 src0{GetSrc<IR::F32>(inst.src[0])};
const IR::F32 src1{GetSrc<IR::F32>(inst.src[1])};
const IR::U1 result = [&] {
@@ -1189,17 +1194,17 @@ void Translator::V_CMP_F32(ConditionOp op, bool set_exec, const GcnInst& inst) {
case ConditionOp::F:
return ir.Imm1(false);
case ConditionOp::EQ:
return ir.FPEqual(src0, src1);
return ir.FPEqual(src0, src1, ordered);
case ConditionOp::LG:
return ir.FPNotEqual(src0, src1);
return ir.FPNotEqual(src0, src1, ordered);
case ConditionOp::GT:
return ir.FPGreaterThan(src0, src1);
return ir.FPGreaterThan(src0, src1, ordered);
case ConditionOp::LT:
return ir.FPLessThan(src0, src1);
return ir.FPLessThan(src0, src1, ordered);
case ConditionOp::LE:
return ir.FPLessThanEqual(src0, src1);
return ir.FPLessThanEqual(src0, src1, ordered);
case ConditionOp::GE:
return ir.FPGreaterThanEqual(src0, src1);
return ir.FPGreaterThanEqual(src0, src1, ordered);
case ConditionOp::U:
return ir.LogicalOr(ir.FPIsNan(src0), ir.FPIsNan(src1));
default:
@@ -1217,7 +1222,7 @@ void Translator::V_CMP_F32(ConditionOp op, bool set_exec, const GcnInst& inst) {
SetDst64(inst.dst[1], ir.Ballot(result));
}
void Translator::V_CMP_F64(ConditionOp op, bool set_exec, const GcnInst& inst) {
void Translator::V_CMP_F64(ConditionOp op, bool ordered, bool set_exec, const GcnInst& inst) {
const IR::F64 src0{GetSrc64<IR::F64>(inst.src[0])};
const IR::F64 src1{GetSrc64<IR::F64>(inst.src[1])};
const IR::U1 result = [&] {
@@ -1225,17 +1230,17 @@ void Translator::V_CMP_F64(ConditionOp op, bool set_exec, const GcnInst& inst) {
case ConditionOp::F:
return ir.Imm1(false);
case ConditionOp::EQ:
return ir.FPEqual(src0, src1);
return ir.FPEqual(src0, src1, ordered);
case ConditionOp::LG:
return ir.FPNotEqual(src0, src1);
return ir.FPNotEqual(src0, src1, ordered);
case ConditionOp::GT:
return ir.FPGreaterThan(src0, src1);
return ir.FPGreaterThan(src0, src1, ordered);
case ConditionOp::LT:
return ir.FPLessThan(src0, src1);
return ir.FPLessThan(src0, src1, ordered);
case ConditionOp::LE:
return ir.FPLessThanEqual(src0, src1);
return ir.FPLessThanEqual(src0, src1, ordered);
case ConditionOp::GE:
return ir.FPGreaterThanEqual(src0, src1);
return ir.FPGreaterThanEqual(src0, src1, ordered);
case ConditionOp::U:
return ir.LogicalOr(ir.FPIsNan(src0), ir.FPIsNan(src1));
default:
@@ -1950,6 +1955,13 @@ void Translator::SetCarryOut(const GcnInst& inst, const IR::U1& carry) {
}
}
IR::F32 Translator::LegacyMul(const IR::F32& a, const IR::F32& b) {
// DX9 rules, 0.0 * x = 0.0
const IR::F32 zero{ir.Imm32(0.0f)};
const IR::U1 either_zero{ir.LogicalOr(ir.FPEqual(a, zero), ir.FPEqual(b, zero))};
return IR::F32{ir.Select(either_zero, zero, ir.FPMul(a, b))};
}
// TODO: add range analysis pass to hopefully put an upper bound on m0, and only select one of
// [src_vgprno, src_vgprno + max_m0]. Same for dst regs we may write back to
src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp
-1
View File
@@ -4,7 +4,6 @@
#include "core/emulator_settings.h"
#include "shader_recompiler/ir/program.h"
#include "shader_recompiler/profile.h"
#include "video_core/buffer_cache/buffer_cache.h"
namespace Shader::Optimization {
src/shader_recompiler/runtime_info.h
+9 -7
View File
@@ -243,13 +243,15 @@ struct RuntimeInfo {
SwStage sw_stage;
HwStage hw_stage;
struct Properties {
AmdGpu::FpDenormMode fp_denorm_mode32;
AmdGpu::FpDenormMode fp_denorm_mode16_64;
AmdGpu::FpRoundMode fp_round_mode32;
AmdGpu::FpRoundMode fp_round_mode16_64;
u32 num_user_data;
u32 num_input_vgprs;
u32 num_allocated_vgprs;
AmdGpu::FpDenormMode fp_denorm_mode32 : 2;
AmdGpu::FpDenormMode fp_denorm_mode16_64 : 2;
AmdGpu::FpRoundMode fp_round_mode32 : 2;
AmdGpu::FpRoundMode fp_round_mode16_64 : 2;
u32 num_allocated_vgprs : 8;
u32 num_user_data : 5;
u32 num_input_vgprs : 2;
u32 dx10_clamp : 1;
u32 reserved : 8;
bool operator==(const Properties&) const noexcept = default;
} props;
union SwInfo {
src/video_core/amdgpu/regs_shader.h
+9 -2
View File
@@ -63,7 +63,10 @@ struct ShaderProgram {
FpRoundMode fp_round_mode64 : 2;
FpDenormMode fp_denorm_mode32 : 2;
FpDenormMode fp_denorm_mode64 : 2;
u32 : 4;
u32 : 1;
u32 dx10_clamp : 1;
u32 debug_mode : 1;
u32 ieee_mode : 1;
u32 vgpr_comp_cnt : 2;
u32 : 6;
u32 scratch_en : 1;
@@ -189,7 +192,11 @@ struct ComputeProgram {
FpRoundMode fp_round_mode64 : 2;
FpDenormMode fp_denorm_mode32 : 2;
FpDenormMode fp_denorm_mode64 : 2;
u32 : 12;
u32 : 1;
u32 dx10_clamp : 1;
u32 debug_mode : 1;
u32 ieee_mode : 1;
u32 : 8;
u32 scratch_en : 1;
u32 num_user_regs : 5;
u32 : 1;
src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
+5 -3
View File
@@ -91,13 +91,14 @@ const Shader::RuntimeInfo& PipelineCache::BuildRuntimeInfo(HwStage stage, SwStag
auto& info = runtime_infos[u32(l_stage)];
const auto& regs = liverpool->regs;
const auto BuildCommon = [&](const auto& program) {
info.props.num_user_data = program.settings.num_user_regs;
info.props.num_input_vgprs = program.settings.vgpr_comp_cnt;
info.props.num_allocated_vgprs = program.NumVgprs();
info.props.fp_denorm_mode32 = program.settings.fp_denorm_mode32;
info.props.fp_denorm_mode16_64 = program.settings.fp_denorm_mode64;
info.props.fp_round_mode32 = program.settings.fp_round_mode32;
info.props.fp_round_mode16_64 = program.settings.fp_round_mode64;
info.props.num_allocated_vgprs = program.NumVgprs();
info.props.num_user_data = program.settings.num_user_regs;
info.props.num_input_vgprs = program.settings.vgpr_comp_cnt;
info.props.dx10_clamp = program.settings.dx10_clamp;
};
info.Initialize(stage, l_stage);
switch (stage) {
@@ -226,6 +227,7 @@ const Shader::RuntimeInfo& PipelineCache::BuildRuntimeInfo(HwStage stage, SwStag
info.props.fp_denorm_mode16_64 = cs_pgm.settings.fp_denorm_mode64;
info.props.fp_round_mode32 = cs_pgm.settings.fp_round_mode32;
info.props.fp_round_mode16_64 = cs_pgm.settings.fp_round_mode64;
info.props.dx10_clamp = cs_pgm.settings.dx10_clamp;
info.hw.cs.workgroup_size = {cs_pgm.num_thread_x.full, cs_pgm.num_thread_y.full,
cs_pgm.num_thread_z.full};
info.hw.cs.tgid_enable = {cs_pgm.IsTgidEnabled(0), cs_pgm.IsTgidEnabled(1),