mirror of
https://github.com/PCSX2/pcsx2.git
synced 2026-10-12 02:08:45 +00:00
GS: Switch CAS from compute shader type to pixel shader type.
Depending on the gpu/driver in some situations it can be a little faster or a little slower. Usually faster on igpus and slower on dgpus, but we want to prioritize igpus since they are on the lower end. GL: Support has been extended to older hardware GL version 3.3. Drop requirement of GL_ARB_gpu_shader5, emulate bitfieldExtract and bitfieldInsert in the cas shader instead. Make GL_ARB_shading_language_packing as a fallback when gl core 4.2 is not available. DX11: Enable on feature level 10. DX12/VK co author port by TJnotJT. Metal co author port by TellowKrinkle.
This commit is contained in:
lightningterror
@@ -640,11 +640,13 @@
|
||||
AU3 AAbsSU3(AU3 a){return AU3(abs(ASU3(a)));}
|
||||
AU4 AAbsSU4(AU4 a){return AU4(abs(ASU4(a)));}
|
||||
//------------------------------------------------------------------------------------------------------------------------------
|
||||
AU1 ABfe(AU1 src,AU1 off,AU1 bits){return bitfieldExtract(src,ASU1(off),ASU1(bits));}
|
||||
// Emulate bitfieldExtract.
|
||||
AU1 ABfe(AU1 src, AU1 off, AU1 bits){uint mask = (bits >= 32u) ? 0xFFFFFFFFu : ((1u << bits) - 1u);return (src >> off) & mask;}
|
||||
AU1 ABfi(AU1 src,AU1 ins,AU1 mask){return (ins&mask)|(src&(~mask));}
|
||||
// Proxy for V_BFI_B32 where the 'mask' is set as 'bits', 'mask=(1<<bits)-1', and 'bits' needs to be an immediate.
|
||||
AU1 ABfiM(AU1 src,AU1 ins,AU1 bits){return bitfieldInsert(src,ins,0,ASU1(bits));}
|
||||
//------------------------------------------------------------------------------------------------------------------------------
|
||||
// Emulate bitfieldInsert.
|
||||
AU1 ABfiM(AU1 src, AU1 ins, AU1 bits){uint mask = (bits >= 32u) ? 0xFFFFFFFFu : ((1u << bits) - 1u);return (src & ~mask) | (ins & mask);}
|
||||
//------------------------------------------------------------------------------------------------------------------------------
|
||||
// V_FRACT_F32 (note DX frac() is different).
|
||||
AF1 AFractF1(AF1 x){return fract(x);}
|
||||
AF2 AFractF2(AF2 x){return fract(x);}
|
||||
|
||||
@@ -19,13 +19,13 @@
|
||||
|
||||
cbuffer cb : register(b0)
|
||||
{
|
||||
uint4 const0;
|
||||
uint4 const1;
|
||||
int2 srcOffset;
|
||||
uint4 const0;
|
||||
uint4 const1;
|
||||
int2 srcOffset;
|
||||
int2 pad;
|
||||
};
|
||||
|
||||
Texture2D InputTexture : register(t0);
|
||||
RWTexture2D<float4> OutputTexture : register(u0);
|
||||
|
||||
#define A_GPU 1
|
||||
#define A_HLSL 1
|
||||
@@ -34,7 +34,7 @@ RWTexture2D<float4> OutputTexture : register(u0);
|
||||
|
||||
AF3 CasLoad(ASU2 p)
|
||||
{
|
||||
return InputTexture.Load(int3(srcOffset, 0) + int3(p, 0)).rgb;
|
||||
return InputTexture.Load(int3(srcOffset, 0) + int3(p, 0)).rgb;
|
||||
}
|
||||
|
||||
// Lets you transform input from the load into a linear color space between 0 and 1. See ffx_cas.h
|
||||
@@ -43,33 +43,19 @@ void CasInput(inout AF1 r, inout AF1 g, inout AF1 b) {}
|
||||
|
||||
#include "ffx_cas.h"
|
||||
|
||||
[numthreads(64, 1, 1)]
|
||||
void main(uint3 LocalThreadId : SV_GroupThreadID, uint3 WorkGroupId : SV_GroupID)
|
||||
float4 ps_main(float4 position : SV_Position, float2 uv : TEXCOORD0) : SV_Target
|
||||
{
|
||||
// Do remapping of local xy in workgroup for a more PS-like swizzle pattern.
|
||||
AU2 gxy = ARmp8x8(LocalThreadId.x) + AU2(WorkGroupId.x << 4u, WorkGroupId.y << 4u);
|
||||
// Pixel coordinate for the current fragment
|
||||
AU2 gxy = AU2(position.xy);
|
||||
|
||||
#if CAS_SHARPEN_ONLY
|
||||
const bool sharpenOnly = true;
|
||||
const bool sharpenOnly = true;
|
||||
#else
|
||||
const bool sharpenOnly = false;
|
||||
const bool sharpenOnly = false;
|
||||
#endif
|
||||
|
||||
// Filter.
|
||||
AF3 c = (float3)0.0f;
|
||||
AF3 c = (AF3)0.0f;
|
||||
CasFilter(c.r, c.g, c.b, gxy, const0, const1, sharpenOnly);
|
||||
|
||||
CasFilter(c.r, c.g, c.b, gxy, const0, const1, sharpenOnly);
|
||||
OutputTexture[ASU2(gxy)] = AF4(c, 1);
|
||||
gxy.x += 8u;
|
||||
|
||||
CasFilter(c.r, c.g, c.b, gxy, const0, const1, sharpenOnly);
|
||||
OutputTexture[ASU2(gxy)] = AF4(c, 1);
|
||||
gxy.y += 8u;
|
||||
|
||||
CasFilter(c.r, c.g, c.b, gxy, const0, const1, sharpenOnly);
|
||||
OutputTexture[ASU2(gxy)] = AF4(c, 1);
|
||||
gxy.x -= 8u;
|
||||
|
||||
CasFilter(c.r, c.g, c.b, gxy, const0, const1, sharpenOnly);
|
||||
OutputTexture[ASU2(gxy)] = AF4(c, 1);
|
||||
}
|
||||
return AF4(c, 1.0f);
|
||||
}
|
||||
@@ -16,12 +16,16 @@
|
||||
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
// THE SOFTWARE.
|
||||
|
||||
//#version 420 // Keep it for editor detection
|
||||
|
||||
#ifdef FRAGMENT_SHADER
|
||||
uniform uvec4 const0;
|
||||
uniform uvec4 const1;
|
||||
uniform ivec2 srcOffset;
|
||||
|
||||
layout(binding=0) uniform sampler2D imgSrc;
|
||||
layout(binding=0, rgba8) uniform writeonly image2D imgDst;
|
||||
layout(binding = 0) uniform sampler2D imgSrc;
|
||||
|
||||
out vec4 fragColor;
|
||||
|
||||
#define A_GPU 1
|
||||
#define A_GLSL 1
|
||||
@@ -39,26 +43,17 @@ void CasInput(inout AF1 r, inout AF1 g, inout AF1 b) {}
|
||||
|
||||
#include "ffx_cas.h"
|
||||
|
||||
layout(local_size_x=64) in;
|
||||
void main()
|
||||
{
|
||||
// Do remapping of local xy in workgroup for a more PS-like swizzle pattern.
|
||||
AU2 gxy = ARmp8x8(gl_LocalInvocationID.x)+AU2(gl_WorkGroupID.x<<4u,gl_WorkGroupID.y<<4u);
|
||||
AU2 gxy = AU2(gl_FragCoord.xy);
|
||||
#if CAS_SHARPEN_ONLY
|
||||
const bool sharpenOnly = true;
|
||||
#else
|
||||
const bool sharpenOnly = false;
|
||||
#endif
|
||||
AF3 c = AF3(0.0f);
|
||||
CasFilter(c.r, c.g, c.b, gxy, const0, const1, sharpenOnly);
|
||||
|
||||
// Filter.
|
||||
AF4 c = vec4(0.0f);
|
||||
CasFilter(c.r, c.g, c.b, gxy, const0, const1, CAS_SHARPEN_ONLY);
|
||||
imageStore(imgDst, ASU2(gxy), c);
|
||||
gxy.x += 8u;
|
||||
|
||||
CasFilter(c.r, c.g, c.b, gxy, const0, const1, CAS_SHARPEN_ONLY);
|
||||
imageStore(imgDst, ASU2(gxy), c);
|
||||
gxy.y += 8u;
|
||||
|
||||
CasFilter(c.r, c.g, c.b, gxy, const0, const1, CAS_SHARPEN_ONLY);
|
||||
imageStore(imgDst, ASU2(gxy), c);
|
||||
gxy.x -= 8u;
|
||||
|
||||
CasFilter(c.r, c.g, c.b, gxy, const0, const1, CAS_SHARPEN_ONLY);
|
||||
imageStore(imgDst, ASU2(gxy), c);
|
||||
fragColor = vec4(c, 1.0f);
|
||||
}
|
||||
#endif
|
||||
@@ -1,6 +1,3 @@
|
||||
#version 460 core
|
||||
#extension GL_EXT_samplerless_texture_functions : require
|
||||
|
||||
// Based on CAS_Shader.glsl
|
||||
// Copyright(c) 2019 Advanced Micro Devices, Inc.All rights reserved.
|
||||
// Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
@@ -21,14 +18,12 @@
|
||||
|
||||
layout(push_constant) uniform const_buffer
|
||||
{
|
||||
uvec4 const0;
|
||||
uvec4 const1;
|
||||
ivec2 srcOffset;
|
||||
uvec4 const0;
|
||||
uvec4 const1;
|
||||
ivec2 srcOffset;
|
||||
};
|
||||
|
||||
layout(set=0, binding=0) uniform texture2D imgSrc;
|
||||
layout(set=0, binding=1, rgba8) uniform writeonly image2D imgDst;
|
||||
layout(constant_id=0) const int sharpenOnly = 0;
|
||||
|
||||
#define A_GPU 1
|
||||
#define A_GLSL 1
|
||||
@@ -37,7 +32,7 @@ layout(constant_id=0) const int sharpenOnly = 0;
|
||||
|
||||
AF3 CasLoad(ASU2 p)
|
||||
{
|
||||
return texelFetch(imgSrc, srcOffset + ivec2(p), 0).rgb;
|
||||
return texelFetch(imgSrc, srcOffset + ivec2(p), 0).rgb;
|
||||
}
|
||||
|
||||
// Lets you transform input from the load into a linear color space between 0 and 1. See ffx_cas.h
|
||||
@@ -46,26 +41,21 @@ void CasInput(inout AF1 r, inout AF1 g, inout AF1 b) {}
|
||||
|
||||
#include "ffx_cas.h"
|
||||
|
||||
layout(local_size_x=64) in;
|
||||
layout(location = 0) out vec4 o_col0;
|
||||
|
||||
void main()
|
||||
{
|
||||
// Do remapping of local xy in workgroup for a more PS-like swizzle pattern.
|
||||
AU2 gxy = ARmp8x8(gl_LocalInvocationID.x)+AU2(gl_WorkGroupID.x<<4u,gl_WorkGroupID.y<<4u);
|
||||
// Pixel coordinate for the current fragment
|
||||
AU2 gxy = AU2(gl_FragCoord.xy);
|
||||
|
||||
// Filter.
|
||||
AF4 c = vec4(0.0f);
|
||||
CasFilter(c.r, c.g, c.b, gxy, const0, const1, sharpenOnly != 0);
|
||||
imageStore(imgDst, ASU2(gxy), c);
|
||||
gxy.x += 8u;
|
||||
#if CAS_SHARPEN_ONLY
|
||||
const bool sharpenOnly = true;
|
||||
#else
|
||||
const bool sharpenOnly = false;
|
||||
#endif
|
||||
|
||||
CasFilter(c.r, c.g, c.b, gxy, const0, const1, sharpenOnly != 0);
|
||||
imageStore(imgDst, ASU2(gxy), c);
|
||||
gxy.y += 8u;
|
||||
AF3 c = AF3(0.0f);
|
||||
CasFilter(c.r, c.g, c.b, gxy, const0, const1, sharpenOnly);
|
||||
|
||||
CasFilter(c.r, c.g, c.b, gxy, const0, const1, sharpenOnly != 0);
|
||||
imageStore(imgDst, ASU2(gxy), c);
|
||||
gxy.x -= 8u;
|
||||
|
||||
CasFilter(c.r, c.g, c.b, gxy, const0, const1, sharpenOnly != 0);
|
||||
imageStore(imgDst, ASU2(gxy), c);
|
||||
}
|
||||
o_col0 = AF4(c, 1.0f);
|
||||
}
|
||||
@@ -1247,10 +1247,10 @@ void GSDevice::CAS(GSTexture*& tex, GSVector4i& src_rect, GSVector4& src_uv, con
|
||||
if (!m_cas || m_cas->GetWidth() != dst_width || m_cas->GetHeight() != dst_height)
|
||||
{
|
||||
delete m_cas;
|
||||
m_cas = CreateSurface(GSTexture::ShaderWriteTexture, dst_width, dst_height, 1, GSTexture::Format::Color);
|
||||
m_cas = CreateSurface(GSTexture::RenderTarget, dst_width, dst_height, 1, GSTexture::Format::Color);
|
||||
if (!m_cas)
|
||||
{
|
||||
Console.Error("Failed to allocate CAS RW texture.");
|
||||
Console.Error("GS: Failed to allocate CAS texture.");
|
||||
return;
|
||||
}
|
||||
}
|
||||
@@ -1265,7 +1265,7 @@ void GSDevice::CAS(GSTexture*& tex, GSVector4i& src_rect, GSVector4& src_uv, con
|
||||
if (!DoCAS(src_tex, m_cas, sharpen_only, consts))
|
||||
{
|
||||
// leave textures intact if we failed
|
||||
Console.Warning("Applying CAS failed.");
|
||||
Console.Warning("GS: Applying CAS failed.");
|
||||
return;
|
||||
}
|
||||
|
||||
|
||||
@@ -592,7 +592,7 @@ bool GSDevice11::Create(GSVSyncMode vsync_mode, bool allow_present_throttle)
|
||||
return false;
|
||||
}
|
||||
|
||||
if (m_features.cas_sharpening && !CreateCASShaders())
|
||||
if (!CreateCASShaders())
|
||||
return false;
|
||||
|
||||
if (!CreateImGuiResources())
|
||||
@@ -715,7 +715,6 @@ void GSDevice11::SetFeatures(IDXGIAdapter1* adapter)
|
||||
SupportsTextureFormat(m_dev.get(), DXGI_FORMAT_BC3_UNORM);
|
||||
|
||||
m_features.bptc_textures = SupportsTextureFormat(m_dev.get(), DXGI_FORMAT_BC7_UNORM);
|
||||
m_features.cas_sharpening = (m_feature_level >= D3D_FEATURE_LEVEL_11_0);
|
||||
m_features.test_and_sample_depth = (m_feature_level >= D3D_FEATURE_LEVEL_11_0);
|
||||
m_features.depth_feedback = m_features.multidraw_fb_copy && GSConfig.DepthFeedbackMode == GSDepthFeedbackMode::Depth;
|
||||
m_features.aa1 = GSConfig.HWAA1 && m_features.vs_expand && m_features.feedback_loops();
|
||||
@@ -2273,7 +2272,7 @@ void GSDevice11::SetupOM(OMDepthStencilSelector dssel, OMBlendSelector bsel, u8
|
||||
bool GSDevice11::CreateCASShaders()
|
||||
{
|
||||
CD3D11_BUFFER_DESC desc(NUM_CAS_CONSTANTS * sizeof(u32), D3D11_BIND_CONSTANT_BUFFER, D3D11_USAGE_DEFAULT);
|
||||
HRESULT hr = m_dev->CreateBuffer(&desc, nullptr, m_cas.cb.put());
|
||||
const HRESULT hr = m_dev->CreateBuffer(&desc, nullptr, m_cas.cb.put());
|
||||
if (FAILED(hr))
|
||||
return false;
|
||||
|
||||
@@ -2281,15 +2280,13 @@ bool GSDevice11::CreateCASShaders()
|
||||
if (!cas_source.has_value() || !GetCASShaderSource(&cas_source.value()))
|
||||
return false;
|
||||
|
||||
static constexpr D3D_SHADER_MACRO sharpen_only_macros[] = {
|
||||
{"CAS_SHARPEN_ONLY", "1"},
|
||||
{nullptr, nullptr}};
|
||||
static constexpr D3D_SHADER_MACRO sharpen_macros[] = {{"CAS_SHARPEN_ONLY", "1"}, {nullptr, nullptr}};
|
||||
|
||||
m_cas.cs_sharpen = m_shader_cache.GetComputeShader(m_dev.get(), cas_source.value(), sharpen_only_macros, "main");
|
||||
m_cas.cs_upscale = m_shader_cache.GetComputeShader(m_dev.get(), cas_source.value(), nullptr, "main");
|
||||
if (!m_cas.cs_sharpen || !m_cas.cs_upscale)
|
||||
m_cas.ps_sharpen = m_shader_cache.GetPixelShader(m_dev.get(), cas_source.value(), sharpen_macros, "ps_main");
|
||||
m_cas.ps_upscale = m_shader_cache.GetPixelShader(m_dev.get(), cas_source.value(), nullptr, "ps_main");
|
||||
if (!m_cas.ps_sharpen || !m_cas.ps_upscale)
|
||||
{
|
||||
Console.Error("D3D11: Failed to create CAS compute shaders.");
|
||||
Console.Error("D3D11: Failed to create CAS pixel shaders.");
|
||||
return false;
|
||||
}
|
||||
|
||||
@@ -2298,27 +2295,14 @@ bool GSDevice11::CreateCASShaders()
|
||||
|
||||
bool GSDevice11::DoCAS(GSTexture* sTex, GSTexture* dTex, bool sharpen_only, const std::array<u32, NUM_CAS_CONSTANTS>& constants)
|
||||
{
|
||||
g_perfmon.Put(GSPerfMon::TextureCopies, 1);
|
||||
const GSVector2i s = dTex->GetSize();
|
||||
|
||||
static constexpr int threadGroupWorkRegionDim = 16;
|
||||
const int dispatchX = (dTex->GetWidth() + (threadGroupWorkRegionDim - 1)) / threadGroupWorkRegionDim;
|
||||
const int dispatchY = (dTex->GetHeight() + (threadGroupWorkRegionDim - 1)) / threadGroupWorkRegionDim;
|
||||
const GSVector4 sRect(0, 0, 1, 1);
|
||||
const GSVector4 dRect(0, 0, s.x, s.y);
|
||||
|
||||
ID3D11ShaderResourceView* srvs[1] = {*static_cast<GSTexture11*>(sTex)};
|
||||
ID3D11UnorderedAccessView* uavs[1] = {*static_cast<GSTexture11*>(dTex)};
|
||||
OMSetRenderTargets(nullptr, nullptr, nullptr);
|
||||
m_ctx->UpdateSubresource(m_cas.cb.get(), 0, nullptr, constants.data(), 0, 0);
|
||||
m_ctx->CSSetConstantBuffers(0, 1, m_cas.cb.addressof());
|
||||
m_ctx->CSSetShader(sharpen_only ? m_cas.cs_sharpen.get() : m_cas.cs_upscale.get(), nullptr, 0);
|
||||
m_ctx->CSSetShaderResources(0, std::size(srvs), srvs);
|
||||
m_ctx->CSSetUnorderedAccessViews(0, std::size(uavs), uavs, nullptr);
|
||||
m_ctx->Dispatch(dispatchX, dispatchY, 1);
|
||||
|
||||
// clear bindings out to prevent hazards
|
||||
uavs[0] = nullptr;
|
||||
srvs[0] = nullptr;
|
||||
m_ctx->CSSetShaderResources(0, std::size(srvs), srvs);
|
||||
m_ctx->CSSetUnorderedAccessViews(0, std::size(uavs), uavs, nullptr);
|
||||
DoStretchRect(sTex, sRect, dTex, dRect, sharpen_only ? m_cas.ps_sharpen.get() : m_cas.ps_upscale.get(), m_cas.cb.get(), Nearest);
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
@@ -271,8 +271,8 @@ private:
|
||||
struct
|
||||
{
|
||||
wil::com_ptr_nothrow<ID3D11Buffer> cb;
|
||||
wil::com_ptr_nothrow<ID3D11ComputeShader> cs_upscale;
|
||||
wil::com_ptr_nothrow<ID3D11ComputeShader> cs_sharpen;
|
||||
wil::com_ptr_nothrow<ID3D11PixelShader> ps_upscale;
|
||||
wil::com_ptr_nothrow<ID3D11PixelShader> ps_sharpen;
|
||||
} m_cas;
|
||||
|
||||
struct
|
||||
|
||||
@@ -60,6 +60,14 @@ static D3D12_RENDER_PASS_BEGINNING_ACCESS_TYPE GetLoadOpForTexture(GSTexture12*
|
||||
// clang-format on
|
||||
}
|
||||
|
||||
static void AddUtilityVertexAttributes(D3D12::GraphicsPipelineBuilder& gpb)
|
||||
{
|
||||
gpb.AddVertexAttribute("POSITION", 0, DXGI_FORMAT_R32G32B32A32_FLOAT, 0, 0);
|
||||
gpb.AddVertexAttribute("TEXCOORD", 0, DXGI_FORMAT_R32G32_FLOAT, 0, 16);
|
||||
gpb.AddVertexAttribute("COLOR", 0, DXGI_FORMAT_R8G8B8A8_UNORM, 0, 28);
|
||||
gpb.SetPrimitiveTopologyType(D3D12_PRIMITIVE_TOPOLOGY_TYPE_TRIANGLE);
|
||||
}
|
||||
|
||||
GSDevice12::ShaderMacro::ShaderMacro()
|
||||
{
|
||||
mlist.emplace_back("DX12", "1");
|
||||
@@ -2327,31 +2335,36 @@ void GSDevice12::DoFXAA(GSTexture* sTex, GSTexture* dTex)
|
||||
|
||||
bool GSDevice12::CompileCASPipelines()
|
||||
{
|
||||
D3D12::RootSignatureBuilder rsb;
|
||||
rsb.Add32BitConstants(0, NUM_CAS_CONSTANTS, D3D12_SHADER_VISIBILITY_ALL);
|
||||
rsb.AddDescriptorTable(D3D12_DESCRIPTOR_RANGE_TYPE_SRV, 0, 1, D3D12_SHADER_VISIBILITY_ALL);
|
||||
rsb.AddDescriptorTable(D3D12_DESCRIPTOR_RANGE_TYPE_UAV, 0, 1, D3D12_SHADER_VISIBILITY_ALL);
|
||||
m_cas_root_signature = rsb.Create(false);
|
||||
if (!m_cas_root_signature)
|
||||
return false;
|
||||
D3D12::GraphicsPipelineBuilder gpb;
|
||||
gpb.SetRootSignature(m_utility_root_signature.get());
|
||||
AddUtilityVertexAttributes(gpb);
|
||||
gpb.SetNoCullRasterizationState();
|
||||
gpb.SetNoBlendingState();
|
||||
gpb.SetVertexShader(m_convert_vs.get());
|
||||
|
||||
std::optional<std::string> cas_source = ReadShaderSource("shaders/dx11/cas.hlsl");
|
||||
if (!cas_source.has_value() || !GetCASShaderSource(&cas_source.value()))
|
||||
return false;
|
||||
|
||||
static constexpr D3D_SHADER_MACRO sharpen_only_macros[] = {{"CAS_SHARPEN_ONLY", "1"}, {nullptr, nullptr}};
|
||||
static constexpr D3D_SHADER_MACRO sharpen_macros[] = {{"CAS_SHARPEN_ONLY", "1"}, {nullptr, nullptr}};
|
||||
|
||||
const ComPtr<ID3DBlob> cs_upscale(m_shader_cache.GetComputeShader(cas_source.value(), nullptr, "main"));
|
||||
const ComPtr<ID3DBlob> cs_sharpen(m_shader_cache.GetComputeShader(cas_source.value(), sharpen_only_macros, "main"));
|
||||
const ComPtr<ID3DBlob> cs_upscale(m_shader_cache.GetPixelShader(cas_source.value(), nullptr, "ps_main"));
|
||||
const ComPtr<ID3DBlob> cs_sharpen(m_shader_cache.GetPixelShader(cas_source.value(), sharpen_macros, "ps_main"));
|
||||
if (!cs_upscale || !cs_sharpen)
|
||||
return false;
|
||||
|
||||
D3D12::ComputePipelineBuilder cpb;
|
||||
cpb.SetRootSignature(m_cas_root_signature.get());
|
||||
cpb.SetShader(cs_upscale->GetBufferPointer(), cs_upscale->GetBufferSize());
|
||||
m_cas_upscale_pipeline = cpb.Create(m_device.get(), m_shader_cache, false);
|
||||
cpb.SetShader(cs_sharpen->GetBufferPointer(), cs_sharpen->GetBufferSize());
|
||||
m_cas_sharpen_pipeline = cpb.Create(m_device.get(), m_shader_cache, false);
|
||||
gpb.SetRootSignature(m_utility_root_signature.get());
|
||||
gpb.SetRenderTarget(0, DXGI_FORMAT_R8G8B8A8_UNORM);
|
||||
gpb.SetNoDepthTestState();
|
||||
gpb.SetNoStencilState();
|
||||
gpb.SetBlendState(0, false, D3D12_BLEND_ONE, D3D12_BLEND_ONE, D3D12_BLEND_OP_ADD, D3D12_BLEND_ZERO,
|
||||
D3D12_BLEND_ZERO, D3D12_BLEND_OP_ADD, D3D12_COLOR_WRITE_ENABLE_ALL);
|
||||
|
||||
gpb.SetPixelShader(cs_upscale.get());
|
||||
m_cas_upscale_pipeline = gpb.Create(m_device.get(), m_shader_cache, false);
|
||||
|
||||
gpb.SetPixelShader(cs_sharpen.get());
|
||||
m_cas_sharpen_pipeline = gpb.Create(m_device.get(), m_shader_cache, false);
|
||||
if (!m_cas_upscale_pipeline || !m_cas_sharpen_pipeline)
|
||||
{
|
||||
Console.Error("D3D12: Failed to create CAS pipelines");
|
||||
@@ -2516,44 +2529,18 @@ void GSDevice12::RenderImGui()
|
||||
bool GSDevice12::DoCAS(
|
||||
GSTexture* sTex, GSTexture* dTex, bool sharpen_only, const std::array<u32, NUM_CAS_CONSTANTS>& constants)
|
||||
{
|
||||
g_perfmon.Put(GSPerfMon::TextureCopies, 1);
|
||||
GL_PUSH("DoCAS");
|
||||
|
||||
EndRenderPass();
|
||||
|
||||
GSTexture12* const sTex12 = static_cast<GSTexture12*>(sTex);
|
||||
GSTexture12* const dTex12 = static_cast<GSTexture12*>(dTex);
|
||||
D3D12DescriptorHandle sTexDH, dTexDH;
|
||||
if (!GetTextureGroupDescriptors(&sTexDH, &sTex12->GetSRVDescriptor(), 1) ||
|
||||
!GetTextureGroupDescriptors(&dTexDH, &dTex12->GetUAVDescriptor(), 1))
|
||||
{
|
||||
ExecuteCommandList(false, "Ran out of descriptors for CAS");
|
||||
if (!GetTextureGroupDescriptors(&sTexDH, &sTex12->GetSRVDescriptor(), 1) ||
|
||||
!GetTextureGroupDescriptors(&dTexDH, &dTex12->GetUAVDescriptor(), 1))
|
||||
{
|
||||
Console.Error("D3D12: Failed to allocate CAS descriptors.");
|
||||
return false;
|
||||
}
|
||||
}
|
||||
SetUtilityRootSignature();
|
||||
SetUtilityPushConstants(constants.data(), sizeof(constants));
|
||||
|
||||
const D3D12CommandList& cmdlist = GetCommandList();
|
||||
const GSTexture12::ResourceState old_state = sTex12->GetResourceState();
|
||||
sTex12->TransitionToState(cmdlist, GSTexture12::ResourceState::ComputeShaderResource);
|
||||
dTex12->TransitionToState(cmdlist, GSTexture12::ResourceState::CASShaderUAV);
|
||||
const GSVector4 dRect(dTex->GetRect());
|
||||
const GSVector4 sRect(0.0f, 0.0f, 1.0f, 1.0f);
|
||||
ID3D12PipelineState* pipeline = sharpen_only ? m_cas_sharpen_pipeline.get() : m_cas_upscale_pipeline.get();
|
||||
DoStretchRect(static_cast<GSTexture12*>(sTex), sRect, static_cast<GSTexture12*>(dTex), dRect, pipeline, Nearest, true);
|
||||
|
||||
cmdlist.list4->SetComputeRootSignature(m_cas_root_signature.get());
|
||||
cmdlist.list4->SetComputeRoot32BitConstants(
|
||||
CAS_ROOT_SIGNATURE_PARAM_PUSH_CONSTANTS, NUM_CAS_CONSTANTS, constants.data(), 0);
|
||||
cmdlist.list4->SetComputeRootDescriptorTable(CAS_ROOT_SIGNATURE_PARAM_SRC_TEXTURE, sTexDH);
|
||||
cmdlist.list4->SetComputeRootDescriptorTable(CAS_ROOT_SIGNATURE_PARAM_DST_TEXTURE, dTexDH);
|
||||
cmdlist.list4->SetPipelineState(sharpen_only ? m_cas_sharpen_pipeline.get() : m_cas_upscale_pipeline.get());
|
||||
m_dirty_flags |= DIRTY_FLAG_PIPELINE;
|
||||
|
||||
static const int threadGroupWorkRegionDim = 16;
|
||||
const int dispatchX = (dTex->GetWidth() + (threadGroupWorkRegionDim - 1)) / threadGroupWorkRegionDim;
|
||||
const int dispatchY = (dTex->GetHeight() + (threadGroupWorkRegionDim - 1)) / threadGroupWorkRegionDim;
|
||||
cmdlist.list4->Dispatch(dispatchX, dispatchY, 1);
|
||||
|
||||
sTex12->TransitionToState(cmdlist, old_state);
|
||||
return true;
|
||||
}
|
||||
|
||||
@@ -2742,14 +2729,6 @@ bool GSDevice12::GetTextureGroupDescriptors(
|
||||
return true;
|
||||
}
|
||||
|
||||
static void AddUtilityVertexAttributes(D3D12::GraphicsPipelineBuilder& gpb)
|
||||
{
|
||||
gpb.AddVertexAttribute("POSITION", 0, DXGI_FORMAT_R32G32B32A32_FLOAT, 0, 0);
|
||||
gpb.AddVertexAttribute("TEXCOORD", 0, DXGI_FORMAT_R32G32_FLOAT, 0, 16);
|
||||
gpb.AddVertexAttribute("COLOR", 0, DXGI_FORMAT_R8G8B8A8_UNORM, 0, 28);
|
||||
gpb.SetPrimitiveTopologyType(D3D12_PRIMITIVE_TOPOLOGY_TYPE_TRIANGLE);
|
||||
}
|
||||
|
||||
GSDevice12::ComPtr<ID3DBlob> GSDevice12::GetUtilityVertexShader(const std::string& source, const char* entry_point)
|
||||
{
|
||||
ShaderMacro sm_model;
|
||||
@@ -2863,14 +2842,17 @@ bool GSDevice12::CreateRootSignatures()
|
||||
|
||||
bool GSDevice12::CompileConvertPipelines()
|
||||
{
|
||||
std::optional<std::string> source = ReadShaderSource("shaders/dx11/convert.fx");
|
||||
if (!source)
|
||||
{
|
||||
Host::ReportErrorAsync("GS", "Failed to read shaders/dx11/convert.fx.");
|
||||
return false;
|
||||
std::optional<std::string> source = ReadShaderSource("shaders/dx11/convert.fx");
|
||||
if (!source)
|
||||
{
|
||||
Host::ReportErrorAsync("GS", "Failed to read shaders/dx11/convert.fx.");
|
||||
return false;
|
||||
}
|
||||
m_convert_source = std::move(*source);
|
||||
}
|
||||
|
||||
m_convert_vs = GetUtilityVertexShader(*source, "vs_main");
|
||||
m_convert_vs = GetUtilityVertexShader(m_convert_source, "vs_main");
|
||||
if (!m_convert_vs)
|
||||
return false;
|
||||
|
||||
@@ -2933,7 +2915,7 @@ bool GSDevice12::CompileConvertPipelines()
|
||||
sm.AddMacro("HAS_FLOAT32_OUTPUT", static_cast<int>(shader.Float32Output()));
|
||||
sm.AddMacro(entry_point_macro.c_str(), 1);
|
||||
|
||||
ComPtr<ID3DBlob> ps(m_shader_cache.GetPixelShader(*source, sm.GetPtr(), shader.EntryPoint()));
|
||||
ComPtr<ID3DBlob> ps(m_shader_cache.GetPixelShader(m_convert_source, sm.GetPtr(), shader.EntryPoint()));
|
||||
if (!ps)
|
||||
return false;
|
||||
|
||||
@@ -2982,7 +2964,7 @@ bool GSDevice12::CompileConvertPipelines()
|
||||
sm.AddMacro("PRIMID_MIN", GSShader::PRIMID_MIN);
|
||||
sm.AddMacro(entry_point_macro.c_str(), "1");
|
||||
|
||||
ComPtr<ID3DBlob> ps(m_shader_cache.GetPixelShader(*source, sm.GetPtr(), entry_point.c_str()));
|
||||
ComPtr<ID3DBlob> ps(m_shader_cache.GetPixelShader(m_convert_source, sm.GetPtr(), entry_point.c_str()));
|
||||
if (!ps)
|
||||
return false;
|
||||
|
||||
|
||||
@@ -445,6 +445,7 @@ private:
|
||||
|
||||
D3D12ShaderCache m_shader_cache;
|
||||
ComPtr<ID3DBlob> m_convert_vs;
|
||||
std::string m_convert_source;
|
||||
std::string m_tfx_source;
|
||||
|
||||
void LookupNativeFormat(GSTexture::Format format, DXGI_FORMAT* d3d_format, DXGI_FORMAT* srv_format,
|
||||
|
||||
@@ -137,7 +137,7 @@ static D3D12_BARRIER_LAYOUT GetD3D12BarrierLayout(GSTexture12::ResourceState sta
|
||||
return D3D12_BARRIER_LAYOUT_DIRECT_QUEUE_COPY_SOURCE;
|
||||
case GSTexture12::ResourceState::CopyDst:
|
||||
return D3D12_BARRIER_LAYOUT_DIRECT_QUEUE_COPY_DEST;
|
||||
case GSTexture12::ResourceState::CASShaderUAV:
|
||||
case GSTexture12::ResourceState::ComputeShaderUAV:
|
||||
case GSTexture12::ResourceState::PixelShaderUAV:
|
||||
return D3D12_BARRIER_LAYOUT_DIRECT_QUEUE_UNORDERED_ACCESS;
|
||||
default:
|
||||
@@ -166,7 +166,7 @@ static D3D12_RESOURCE_STATES GetD3D12ResourceState(GSTexture12::ResourceState st
|
||||
return D3D12_RESOURCE_STATE_COPY_SOURCE;
|
||||
case GSTexture12::ResourceState::CopyDst:
|
||||
return D3D12_RESOURCE_STATE_COPY_DEST;
|
||||
case GSTexture12::ResourceState::CASShaderUAV:
|
||||
case GSTexture12::ResourceState::ComputeShaderUAV:
|
||||
case GSTexture12::ResourceState::PixelShaderUAV:
|
||||
return D3D12_RESOURCE_STATE_UNORDERED_ACCESS;
|
||||
default:
|
||||
@@ -876,7 +876,7 @@ void GSTexture12::TransitionSubresourceToState(const D3D12CommandList& cmdlist,
|
||||
barrier.AccessBefore = D3D12_BARRIER_ACCESS_COPY_DEST;
|
||||
barrier.SyncBefore = D3D12_BARRIER_SYNC_COPY;
|
||||
break;
|
||||
case ResourceState::CASShaderUAV:
|
||||
case ResourceState::ComputeShaderUAV:
|
||||
barrier.LayoutBefore = m_simultaneous_tex ? D3D12_BARRIER_LAYOUT_COMMON : D3D12_BARRIER_LAYOUT_DIRECT_QUEUE_UNORDERED_ACCESS;
|
||||
barrier.AccessBefore = D3D12_BARRIER_ACCESS_UNORDERED_ACCESS;
|
||||
barrier.SyncBefore = D3D12_BARRIER_SYNC_COMPUTE_SHADING;
|
||||
@@ -954,7 +954,7 @@ void GSTexture12::TransitionSubresourceToState(const D3D12CommandList& cmdlist,
|
||||
barrier.AccessAfter = D3D12_BARRIER_ACCESS_COPY_DEST;
|
||||
barrier.SyncAfter = D3D12_BARRIER_SYNC_COPY;
|
||||
break;
|
||||
case ResourceState::CASShaderUAV:
|
||||
case ResourceState::ComputeShaderUAV:
|
||||
barrier.LayoutAfter = m_simultaneous_tex ? D3D12_BARRIER_LAYOUT_COMMON : D3D12_BARRIER_LAYOUT_DIRECT_QUEUE_UNORDERED_ACCESS;
|
||||
barrier.AccessAfter = D3D12_BARRIER_ACCESS_UNORDERED_ACCESS;
|
||||
barrier.SyncAfter = D3D12_BARRIER_SYNC_COMPUTE_SHADING;
|
||||
@@ -1042,10 +1042,10 @@ void GSTexture12::TransitionSubresourceToState(const D3D12CommandList& cmdlist,
|
||||
case ResourceState::CopyDst:
|
||||
barrier.Transition.StateBefore = D3D12_RESOURCE_STATE_COPY_DEST;
|
||||
break;
|
||||
case ResourceState::CASShaderUAV:
|
||||
case ResourceState::ComputeShaderUAV:
|
||||
case ResourceState::PixelShaderUAV:
|
||||
// Handled in after_state cases.
|
||||
if (after_state == ResourceState::CASShaderUAV || after_state == ResourceState::PixelShaderUAV)
|
||||
if (after_state == ResourceState::ComputeShaderUAV || after_state == ResourceState::PixelShaderUAV)
|
||||
break;
|
||||
|
||||
barrier.Transition.StateBefore = D3D12_RESOURCE_STATE_UNORDERED_ACCESS;
|
||||
@@ -1093,9 +1093,9 @@ void GSTexture12::TransitionSubresourceToState(const D3D12CommandList& cmdlist,
|
||||
case ResourceState::CopyDst:
|
||||
barrier.Transition.StateAfter = D3D12_RESOURCE_STATE_COPY_DEST;
|
||||
break;
|
||||
case ResourceState::CASShaderUAV:
|
||||
case ResourceState::ComputeShaderUAV:
|
||||
case ResourceState::PixelShaderUAV:
|
||||
if (before_state == ResourceState::CASShaderUAV || before_state == ResourceState::PixelShaderUAV)
|
||||
if (before_state == ResourceState::ComputeShaderUAV || before_state == ResourceState::PixelShaderUAV)
|
||||
{
|
||||
// No state transition, but probably want a barrier instead.
|
||||
barrier.Type = D3D12_RESOURCE_BARRIER_TYPE_UAV;
|
||||
|
||||
@@ -32,7 +32,7 @@ public:
|
||||
ComputeShaderResource,
|
||||
CopySrc,
|
||||
CopyDst,
|
||||
CASShaderUAV, // No Clear UAV Sync
|
||||
ComputeShaderUAV, // No Clear UAV Sync
|
||||
PixelShaderUAV,
|
||||
Count
|
||||
};
|
||||
|
||||
@@ -248,7 +248,7 @@ public:
|
||||
MRCOwned<id<MTLFence>> m_spin_fence;
|
||||
|
||||
// Functions and Pipeline States
|
||||
MRCOwned<id<MTLComputePipelineState>> m_cas_pipeline[2];
|
||||
MRCOwned<id<MTLRenderPipelineState>> m_cas_pipeline[2];
|
||||
std::vector<MRCOwned<id<MTLRenderPipelineState>>> m_convert_pipeline;
|
||||
MRCOwned<id<MTLRenderPipelineState>> m_present_pipeline[static_cast<int>(PresentShader::Count)];
|
||||
MRCOwned<id<MTLRenderPipelineState>> m_merge_pipeline[4];
|
||||
|
||||
@@ -749,23 +749,12 @@ void GSDeviceMTL::DoShadeBoost(GSTexture* sTex, GSTexture* dTex, const float par
|
||||
|
||||
bool GSDeviceMTL::DoCAS(GSTexture* sTex, GSTexture* dTex, bool sharpen_only, const std::array<u32, NUM_CAS_CONSTANTS>& constants)
|
||||
{ @autoreleasepool {
|
||||
g_perfmon.Put(GSPerfMon::TextureCopies, 1);
|
||||
|
||||
static constexpr int threadGroupWorkRegionDim = 16;
|
||||
const int dispatchX = (dTex->GetWidth() + (threadGroupWorkRegionDim - 1)) / threadGroupWorkRegionDim;
|
||||
const int dispatchY = (dTex->GetHeight() + (threadGroupWorkRegionDim - 1)) / threadGroupWorkRegionDim;
|
||||
static_assert(sizeof(constants) == sizeof(GSMTLCASPSUniform));
|
||||
|
||||
EndRenderPass();
|
||||
id<MTLComputeCommandEncoder> enc = [GetRenderCmdBuf() computeCommandEncoder];
|
||||
[enc setLabel:@"CAS"];
|
||||
[enc setComputePipelineState:m_cas_pipeline[sharpen_only]];
|
||||
[enc setTexture:static_cast<GSTextureMTL*>(sTex)->GetTexture() atIndex:0];
|
||||
[enc setTexture:static_cast<GSTextureMTL*>(dTex)->GetTexture() atIndex:1];
|
||||
[enc setBytes:&constants length:sizeof(constants) atIndex:GSMTLBufferIndexUniforms];
|
||||
[enc dispatchThreadgroups:MTLSizeMake(dispatchX, dispatchY, 1)
|
||||
threadsPerThreadgroup:MTLSizeMake(64, 1, 1)];
|
||||
[enc endEncoding];
|
||||
BeginRenderPass(@"CAS", dTex, MTLLoadActionDontCare, nullptr, MTLLoadActionDontCare);
|
||||
[m_current_render.encoder setFragmentBytes:&constants
|
||||
length:sizeof(constants)
|
||||
atIndex:GSMTLBufferIndexUniforms];
|
||||
RenderCopy(sTex, m_cas_pipeline[sharpen_only], GSVector4i(0, 0, dTex->GetSize().x, dTex->GetSize().y));
|
||||
return true;
|
||||
}}
|
||||
|
||||
@@ -1095,13 +1084,6 @@ bool GSDeviceMTL::Create(GSVSyncMode vsync_mode, bool allow_present_throttle)
|
||||
[clearSpinBuffer endEncoding];
|
||||
m_spin_pipeline = MakeComputePipeline(LoadShader(@"waste_time"), @"waste_time");
|
||||
|
||||
for (int sharpen_only = 0; sharpen_only < 2; sharpen_only++)
|
||||
{
|
||||
setFnConstantB(m_fn_constants, sharpen_only, GSMTLConstantIndex_CAS_SHARPEN_ONLY);
|
||||
NSString* shader = m_dev.features.has_fast_half ? @"CASHalf" : @"CASFloat";
|
||||
m_cas_pipeline[sharpen_only] = MakeComputePipeline(LoadShader(shader), sharpen_only ? @"CAS Sharpen" : @"CAS Upscale");
|
||||
}
|
||||
|
||||
m_expand_index_buffer = CreatePrivateBufferWithContent(m_dev.dev, initCommands, MTLResourceHazardTrackingModeUntracked, EXPAND_BUFFER_SIZE, GenerateExpansionIndexBuffer);
|
||||
[m_expand_index_buffer setLabel:@"Point/Sprite Expand Indices"];
|
||||
|
||||
@@ -1262,6 +1244,13 @@ bool GSDeviceMTL::Create(GSVSyncMode vsync_mode, bool allow_present_throttle)
|
||||
m_primid_init_pipeline[0][3] = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_primid_rta_init_datm1"), @"PrimID DATM1 RTA Clear");
|
||||
|
||||
pdesc.colorAttachments[0].pixelFormat = ConvertPixelFormat(GSTexture::Format::Color);
|
||||
|
||||
for (int sharpen_only = 0; sharpen_only < 2; sharpen_only++)
|
||||
{
|
||||
setFnConstantB(m_fn_constants, sharpen_only, GSMTLConstantIndex_CAS_SHARPEN_ONLY);
|
||||
m_cas_pipeline[sharpen_only] = MakePipeline(pdesc, fs_triangle, LoadShader(@"CASPS"), sharpen_only ? @"CAS Sharpen" : @"CAS Upscale");
|
||||
}
|
||||
|
||||
applyAttribute(pdesc.vertexDescriptor, 0, MTLVertexFormatFloat2, offsetof(ConvertShaderVertex, pos), 0);
|
||||
applyAttribute(pdesc.vertexDescriptor, 1, MTLVertexFormatFloat2, offsetof(ConvertShaderVertex, texpos), 0);
|
||||
pdesc.vertexDescriptor.layouts[0].stride = sizeof(ConvertShaderVertex);
|
||||
|
||||
@@ -41,6 +41,19 @@ A_STATIC void CasInputH(inoutAH2 r, inoutAH2 g, inoutAH2 b){}
|
||||
|
||||
constant bool CAS_SHARPEN_ONLY [[function_constant(GSMTLConstantIndex_CAS_SHARPEN_ONLY)]];
|
||||
|
||||
fragment float4 CASPS(
|
||||
float4 pos [[position]],
|
||||
texture2d<float, access::read> input [[texture(GSMTLTextureIndexNonHW)]],
|
||||
constant GSMTLCASPSUniform& cb [[buffer(GSMTLBufferIndexUniforms)]])
|
||||
{
|
||||
const CASTextureF tex{input, AU2(cb.srcOffset)};
|
||||
const AU4 const0 = cb.const0;
|
||||
const AU4 const1 = cb.const1;
|
||||
float r, g, b;
|
||||
CasFilter(tex, r, g, b, uint2(pos.xy), const0, const1, CAS_SHARPEN_ONLY);
|
||||
return float4(r, g, b, 1);
|
||||
}
|
||||
|
||||
kernel void CASFloat(
|
||||
uint2 localID [[thread_position_in_threadgroup]],
|
||||
uint2 workgroupID [[threadgroup_position_in_grid]],
|
||||
|
||||
@@ -653,8 +653,8 @@ bool GSDeviceOGL::Create(GSVSyncMode vsync_mode, bool allow_present_throttle)
|
||||
if (!CompileShadeBoostProgram() || !CompileFXAAProgram())
|
||||
return false;
|
||||
|
||||
// Image load store and GLSL 420pack is core in GL4.2, no need to check.
|
||||
m_features.cas_sharpening = ((GLAD_GL_VERSION_4_2 && GLAD_GL_ARB_compute_shader) || GLAD_GL_ES_VERSION_3_2) && CreateCASPrograms();
|
||||
if (GLAD_GL_VERSION_4_2 || GLAD_GL_ARB_shading_language_packing)
|
||||
m_features.cas_sharpening = CreateCASPrograms();
|
||||
|
||||
// ****************************************************************
|
||||
// rasterization configuration
|
||||
@@ -1020,6 +1020,19 @@ bool GSDeviceOGL::CheckFeatures()
|
||||
GLAD_GL_AMD_conservative_depth ? "GL_AMD_conservative_depth" :
|
||||
"None");
|
||||
|
||||
Console.WriteLnFmt(
|
||||
"GL: {1:<{0}} {3:<{2}} => Using: {4}",
|
||||
LABEL_WIDTH, "CAS Sharpening:",
|
||||
STATUS_WIDTH,
|
||||
(GLAD_GL_VERSION_4_2 || GLAD_GL_ARB_shading_language_packing) ?
|
||||
"Supported" :
|
||||
"Not Supported",
|
||||
GLAD_GL_VERSION_4_2 ?
|
||||
"OpenGL 4.2 Core" :
|
||||
GLAD_GL_ARB_shading_language_packing ?
|
||||
"GL_ARB_shading_language_packing" :
|
||||
"None");
|
||||
|
||||
Console.WriteLnFmt(
|
||||
"GL: {1:<{0}} {3:<{2}} => Using: {4}",
|
||||
LABEL_WIDTH, "Texture Barriers:",
|
||||
@@ -1776,6 +1789,11 @@ std::string GSDeviceOGL::GenGlslHeader(const std::string_view entry, GLenum type
|
||||
else
|
||||
header += "#define HAS_CLIP_CONTROL 0\n";
|
||||
|
||||
if (!GLAD_GL_VERSION_4_2 && GLAD_GL_ARB_shading_language_packing)
|
||||
{
|
||||
header += "#extension GL_ARB_shading_language_packing : require\n";
|
||||
}
|
||||
|
||||
// Allow to puts several shader in 1 files
|
||||
switch (type)
|
||||
{
|
||||
@@ -2583,25 +2601,28 @@ void GSDeviceOGL::ClearSamplerCache()
|
||||
|
||||
bool GSDeviceOGL::CreateCASPrograms()
|
||||
{
|
||||
std::optional<std::string> cas_source = ReadShaderSource("shaders/opengl/cas.glsl");
|
||||
if (!cas_source.has_value() || !GetCASShaderSource(&cas_source.value()))
|
||||
std::optional<std::string> shader = ReadShaderSource("shaders/opengl/cas.glsl");
|
||||
if (!shader.has_value() || !GetCASShaderSource(&shader.value()))
|
||||
{
|
||||
Console.Error("GL: Failed to read cas.glsl");
|
||||
m_features.cas_sharpening = false;
|
||||
return false;
|
||||
}
|
||||
|
||||
const char* header =
|
||||
"#version 420\n"
|
||||
"#extension GL_ARB_compute_shader : require\n";
|
||||
const char* sharpen_params[2] = {
|
||||
"#define CAS_SHARPEN_ONLY false\n",
|
||||
"#define CAS_SHARPEN_ONLY true\n"};
|
||||
const std::array<std::pair<GLProgram*, const char*>, 2> programs = {{
|
||||
{&m_cas.upscale_ps, "#define CAS_SHARPEN_ONLY 0\n"},
|
||||
{&m_cas.sharpen_ps, "#define CAS_SHARPEN_ONLY 1\n"},
|
||||
}};
|
||||
|
||||
if (!m_shader_cache.GetComputeProgram(&m_cas.upscale_ps, fmt::format("{}{}{}", header, sharpen_params[0], cas_source.value())) ||
|
||||
!m_shader_cache.GetComputeProgram(&m_cas.sharpen_ps, fmt::format("{}{}{}", header, sharpen_params[1], cas_source.value())))
|
||||
for (const auto& [prog, macro] : programs)
|
||||
{
|
||||
m_features.cas_sharpening = false;
|
||||
return false;
|
||||
const std::string ps(GetShaderSource("main", GL_FRAGMENT_SHADER, *shader, macro));
|
||||
if (!m_shader_cache.GetProgram(prog, m_convert.vs, ps))
|
||||
{
|
||||
Console.Error("GL: Failed to compile CAS program.");
|
||||
m_features.cas_sharpening = false;
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
const auto link_uniforms = [](GLProgram& prog) {
|
||||
@@ -2609,6 +2630,7 @@ bool GSDeviceOGL::CreateCASPrograms()
|
||||
prog.RegisterUniform("const1");
|
||||
prog.RegisterUniform("srcOffset");
|
||||
};
|
||||
|
||||
link_uniforms(m_cas.upscale_ps);
|
||||
link_uniforms(m_cas.sharpen_ps);
|
||||
|
||||
@@ -2617,21 +2639,22 @@ bool GSDeviceOGL::CreateCASPrograms()
|
||||
|
||||
bool GSDeviceOGL::DoCAS(GSTexture* sTex, GSTexture* dTex, bool sharpen_only, const std::array<u32, NUM_CAS_CONSTANTS>& constants)
|
||||
{
|
||||
g_perfmon.Put(GSPerfMon::TextureCopies, 1);
|
||||
GL_PUSH("DoCAS");
|
||||
|
||||
OMSetColorMaskState();
|
||||
|
||||
const GSVector2i s = dTex->GetSize();
|
||||
const GSVector4 sRect(0, 0, 1, 1);
|
||||
const GSVector4 dRect(0, 0, s.x, s.y);
|
||||
|
||||
const GLProgram& prog = sharpen_only ? m_cas.sharpen_ps : m_cas.upscale_ps;
|
||||
prog.Bind();
|
||||
prog.Uniform4uiv(0, &constants[0]);
|
||||
prog.Uniform4uiv(1, &constants[4]);
|
||||
prog.Uniform2iv(2, reinterpret_cast<const s32*>(&constants[8]));
|
||||
|
||||
PSSetShaderResource(TEXTURE_TEXTURE, sTex);
|
||||
glBindImageTexture(0, static_cast<GSTextureOGL*>(dTex)->GetID(), 0, GL_FALSE, 0, GL_WRITE_ONLY, GL_RGBA8);
|
||||
prog.Uniform4uiv(0, &constants[0]); // const0
|
||||
prog.Uniform4uiv(1, &constants[4]); // const1
|
||||
prog.Uniform2iv(2, reinterpret_cast<const s32*>(&constants[8])); // srcOffset
|
||||
|
||||
static const int threadGroupWorkRegionDim = 16;
|
||||
const int dispatchX = (dTex->GetWidth() + (threadGroupWorkRegionDim - 1)) / threadGroupWorkRegionDim;
|
||||
const int dispatchY = (dTex->GetHeight() + (threadGroupWorkRegionDim - 1)) / threadGroupWorkRegionDim;
|
||||
glDispatchCompute(dispatchX, dispatchY, 1);
|
||||
DoStretchRect(sTex, sRect, dTex, dRect, prog, Nearest);
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
@@ -4154,14 +4154,17 @@ bool GSDeviceVK::CreateRenderPasses()
|
||||
|
||||
bool GSDeviceVK::CompileConvertPipelines()
|
||||
{
|
||||
const std::optional<std::string> source = ReadShaderSource("shaders/vulkan/convert.glsl");
|
||||
if (!source)
|
||||
{
|
||||
Host::ReportErrorAsync("GS", "Failed to read shaders/vulkan/convert.glsl.");
|
||||
return false;
|
||||
const std::optional<std::string> source = ReadShaderSource("shaders/vulkan/convert.glsl");
|
||||
if (!source)
|
||||
{
|
||||
Host::ReportErrorAsync("GS", "Failed to read shaders/vulkan/convert.glsl.");
|
||||
return false;
|
||||
}
|
||||
m_convert_source = std::move(*source);
|
||||
}
|
||||
|
||||
VkShaderModule vs = GetUtilityVertexShader(*source);
|
||||
VkShaderModule vs = GetUtilityVertexShader(m_convert_source);
|
||||
if (vs == VK_NULL_HANDLE)
|
||||
return false;
|
||||
ScopedGuard vs_guard([this, &vs]() { vkDestroyShaderModule(m_device, vs, nullptr); });
|
||||
@@ -4232,7 +4235,7 @@ bool GSDeviceVK::CompileConvertPipelines()
|
||||
macro += fmt::format("#define HAS_FLOAT32_INPUT {}\n", static_cast<int>(shader.Float32Input()));
|
||||
macro += fmt::format("#define HAS_FLOAT32_OUTPUT {}\n", static_cast<int>(shader.Float32Output()));
|
||||
|
||||
std::string shader_with_header = macro + *source;
|
||||
std::string shader_with_header = macro + m_convert_source;
|
||||
|
||||
VkShaderModule ps = GetUtilityFragmentShader(shader_with_header, shader.EntryPoint());
|
||||
if (ps == VK_NULL_HANDLE)
|
||||
@@ -4296,7 +4299,7 @@ bool GSDeviceVK::CompileConvertPipelines()
|
||||
macro += fmt::format("#define PRIMID_MAX {}\n", GSShader::PRIMID_MAX);
|
||||
macro += fmt::format("#define PRIMID_MIN {}\n", GSShader::PRIMID_MIN);
|
||||
|
||||
const std::string source_with_header = macro + *source;
|
||||
const std::string source_with_header = macro + m_convert_source;
|
||||
|
||||
const std::string entry_point(StringUtil::StdStringFromFormat("ps_primid_image_init_%d", datm));
|
||||
VkShaderModule ps = GetUtilityFragmentShader(source_with_header, entry_point.c_str());
|
||||
@@ -4574,40 +4577,47 @@ bool GSDeviceVK::CompilePostProcessingPipelines()
|
||||
|
||||
bool GSDeviceVK::CompileCASPipelines()
|
||||
{
|
||||
VkShaderModule vs = GetUtilityVertexShader(m_convert_source);
|
||||
if (vs == VK_NULL_HANDLE)
|
||||
return false;
|
||||
ScopedGuard vs_guard([this, &vs]() { vkDestroyShaderModule(m_device, vs, nullptr); });
|
||||
|
||||
VkDevice dev = m_device;
|
||||
Vulkan::DescriptorSetLayoutBuilder dslb;
|
||||
Vulkan::PipelineLayoutBuilder plb;
|
||||
|
||||
dslb.SetPushFlag();
|
||||
dslb.AddBinding(0, VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE, 1, VK_SHADER_STAGE_COMPUTE_BIT);
|
||||
dslb.AddBinding(1, VK_DESCRIPTOR_TYPE_STORAGE_IMAGE, 1, VK_SHADER_STAGE_COMPUTE_BIT);
|
||||
if ((m_cas_ds_layout = dslb.Create(dev)) == VK_NULL_HANDLE)
|
||||
return false;
|
||||
Vulkan::SetObjectName(dev, m_cas_ds_layout, "CAS descriptor layout");
|
||||
Vulkan::GraphicsPipelineBuilder gpb;
|
||||
SetPipelineProvokingVertex(m_features, gpb);
|
||||
AddUtilityVertexAttributes(gpb);
|
||||
gpb.SetPipelineLayout(m_utility_pipeline_layout);
|
||||
gpb.SetDynamicViewportAndScissorState();
|
||||
gpb.AddDynamicState(VK_DYNAMIC_STATE_BLEND_CONSTANTS);
|
||||
gpb.AddDynamicState(VK_DYNAMIC_STATE_LINE_WIDTH);
|
||||
gpb.SetNoCullRasterizationState();
|
||||
gpb.SetNoBlendingState();
|
||||
gpb.SetVertexShader(vs);
|
||||
|
||||
plb.AddPushConstants(VK_SHADER_STAGE_COMPUTE_BIT, 0, NUM_CAS_CONSTANTS * sizeof(u32));
|
||||
plb.AddDescriptorSet(m_cas_ds_layout);
|
||||
if ((m_cas_pipeline_layout = plb.Create(dev)) == VK_NULL_HANDLE)
|
||||
return false;
|
||||
Vulkan::SetObjectName(dev, m_cas_pipeline_layout, "CAS pipeline layout");
|
||||
gpb.SetRenderPass(
|
||||
GetRenderPass(VK_FORMAT_R8G8B8A8_UNORM, VK_FORMAT_UNDEFINED, VK_ATTACHMENT_LOAD_OP_DONT_CARE), 0);
|
||||
|
||||
// we use specialization constants to avoid compiling it twice
|
||||
std::optional<std::string> cas_source = ReadShaderSource("shaders/vulkan/cas.glsl");
|
||||
if (!cas_source.has_value() || !GetCASShaderSource(&cas_source.value()))
|
||||
return false;
|
||||
|
||||
VkShaderModule mod = g_vulkan_shader_cache->GetComputeShader(cas_source->c_str());
|
||||
ScopedGuard mod_guard = [this, &mod]() { vkDestroyShaderModule(m_device, mod, nullptr); };
|
||||
if (mod == VK_NULL_HANDLE)
|
||||
return false;
|
||||
|
||||
for (u8 sharpen_only = 0; sharpen_only < 2; sharpen_only++)
|
||||
{
|
||||
Vulkan::ComputePipelineBuilder cpb;
|
||||
cpb.SetPipelineLayout(m_cas_pipeline_layout);
|
||||
cpb.SetShader(mod, "main");
|
||||
cpb.SetSpecializationBool(0, sharpen_only != 0);
|
||||
m_cas_pipelines[sharpen_only] = cpb.Create(dev, g_vulkan_shader_cache->GetPipelineCache(true), false);
|
||||
std::stringstream source_with_header;
|
||||
AddMacro(source_with_header, "CAS_SHARPEN_ONLY", sharpen_only);
|
||||
|
||||
source_with_header << *cas_source;
|
||||
|
||||
VkShaderModule mod = GetUtilityFragmentShader(source_with_header.str());
|
||||
|
||||
ScopedGuard mod_guard = [this, &mod]() { vkDestroyShaderModule(m_device, mod, nullptr); };
|
||||
if (mod == VK_NULL_HANDLE)
|
||||
return false;
|
||||
|
||||
gpb.SetFragmentShader(mod);
|
||||
|
||||
m_cas_pipelines[sharpen_only] = gpb.Create(dev, g_vulkan_shader_cache->GetPipelineCache(true), false);
|
||||
if (!m_cas_pipelines[sharpen_only])
|
||||
return false;
|
||||
}
|
||||
@@ -4771,7 +4781,7 @@ void GSDeviceVK::RenderBlankFrame()
|
||||
bool GSDeviceVK::DoCAS(
|
||||
GSTexture* sTex, GSTexture* dTex, bool sharpen_only, const std::array<u32, NUM_CAS_CONSTANTS>& constants)
|
||||
{
|
||||
g_perfmon.Put(GSPerfMon::TextureCopies, 1);
|
||||
GL_PUSH("DoCAS");
|
||||
|
||||
EndRenderPass();
|
||||
|
||||
@@ -4780,27 +4790,17 @@ bool GSDeviceVK::DoCAS(
|
||||
VkCommandBuffer cmdbuf = GetCurrentCommandBuffer();
|
||||
|
||||
sTexVK->TransitionToLayout(cmdbuf, GSTextureVK::Layout::ShaderReadOnly);
|
||||
dTexVK->TransitionToLayout(cmdbuf, GSTextureVK::Layout::ComputeReadWriteImage);
|
||||
dTexVK->TransitionToLayout(cmdbuf, GSTextureVK::Layout::ColorAttachment);
|
||||
|
||||
// only happening once a frame, so the update isn't a huge deal.
|
||||
Vulkan::DescriptorSetUpdateBuilder dsub;
|
||||
dsub.AddImageDescriptorWrite(VK_NULL_HANDLE, 0, sTexVK->GetView(), sTexVK->GetVkLayout());
|
||||
dsub.AddStorageImageDescriptorWrite(VK_NULL_HANDLE, 1, dTexVK->GetView(), dTexVK->GetVkLayout());
|
||||
dsub.PushUpdate(cmdbuf, VK_PIPELINE_BIND_POINT_COMPUTE, m_cas_pipeline_layout, 0, false);
|
||||
SetUtilityPushConstants(constants.data(), NUM_CAS_CONSTANTS * sizeof(u32));
|
||||
const GSVector4 dRect(dTex->GetRect());
|
||||
const GSVector4 sRect(0.0f, 0.0f, 1.0f, 1.0f);
|
||||
const VkPipeline pipeline = m_cas_pipelines[static_cast<u8>(sharpen_only)];
|
||||
DoStretchRect(static_cast<GSTextureVK*>(sTex), sRect, static_cast<GSTextureVK*>(dTex), dRect, pipeline, Nearest, true);
|
||||
|
||||
// the actual meat and potatoes! only four commands.
|
||||
static const int threadGroupWorkRegionDim = 16;
|
||||
const int dispatchX = (dTex->GetWidth() + (threadGroupWorkRegionDim - 1)) / threadGroupWorkRegionDim;
|
||||
const int dispatchY = (dTex->GetHeight() + (threadGroupWorkRegionDim - 1)) / threadGroupWorkRegionDim;
|
||||
EndRenderPass();
|
||||
static_cast<GSTextureVK*>(dTex)->TransitionToLayout(GSTextureVK::Layout::ShaderReadOnly);
|
||||
|
||||
vkCmdPushConstants(cmdbuf, m_cas_pipeline_layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, NUM_CAS_CONSTANTS * sizeof(u32),
|
||||
constants.data());
|
||||
vkCmdBindPipeline(cmdbuf, VK_PIPELINE_BIND_POINT_COMPUTE, m_cas_pipelines[static_cast<u8>(sharpen_only)]);
|
||||
vkCmdDispatch(cmdbuf, dispatchX, dispatchY, 1);
|
||||
|
||||
dTexVK->TransitionToLayout(GSTextureVK::Layout::ShaderReadOnly);
|
||||
|
||||
// all done!
|
||||
return true;
|
||||
}
|
||||
|
||||
@@ -4863,10 +4863,7 @@ void GSDeviceVK::DestroyResources()
|
||||
if (it != VK_NULL_HANDLE)
|
||||
vkDestroyPipeline(m_device, it, nullptr);
|
||||
}
|
||||
if (m_cas_pipeline_layout != VK_NULL_HANDLE)
|
||||
vkDestroyPipelineLayout(m_device, m_cas_pipeline_layout, nullptr);
|
||||
if (m_cas_ds_layout != VK_NULL_HANDLE)
|
||||
vkDestroyDescriptorSetLayout(m_device, m_cas_ds_layout, nullptr);
|
||||
|
||||
if (m_imgui_pipeline != VK_NULL_HANDLE)
|
||||
vkDestroyPipeline(m_device, m_imgui_pipeline, nullptr);
|
||||
|
||||
|
||||
@@ -472,8 +472,6 @@ private:
|
||||
|
||||
VkRenderPass m_tfx_render_pass[2][2][2][3][2][2][3][3] = {}; // [rt][ds][colclip][date][fbl][dsp][rt_op][ds_op]
|
||||
|
||||
VkDescriptorSetLayout m_cas_ds_layout = VK_NULL_HANDLE;
|
||||
VkPipelineLayout m_cas_pipeline_layout = VK_NULL_HANDLE;
|
||||
std::array<VkPipeline, NUM_CAS_PIPELINES> m_cas_pipelines = {};
|
||||
VkPipeline m_imgui_pipeline = VK_NULL_HANDLE;
|
||||
|
||||
@@ -481,6 +479,7 @@ private:
|
||||
GSHWDrawConfig::PSConstantBuffer m_ps_cb_cache;
|
||||
GSHWDrawConfig::VSPushConstants m_vs_pc_cache;
|
||||
|
||||
std::string m_convert_source;
|
||||
std::string m_tfx_source;
|
||||
|
||||
GSTexture* CreateSurface(GSTexture::Usage usage, int width, int height, int levels, GSTexture::Format format) override;
|
||||
|
||||
@@ -3,4 +3,4 @@
|
||||
|
||||
/// Version number for GS and other shaders. Increment whenever any of the contents of the
|
||||
/// shaders change, to invalidate the cache.
|
||||
static constexpr u32 SHADER_CACHE_VERSION = 115; // Last changed in PR 14985
|
||||
static constexpr u32 SHADER_CACHE_VERSION = 116; // Last changed in PR 14992
|
||||
|
||||
Reference in New Issue
Block a user