diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 850ec7c8..874e8b90 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -66,6 +66,7 @@ if (ENABLE_OGLRENDERER) GPU_OpenGL.cpp GPU_OpenGL_shaders.h GPU3D_OpenGL.cpp + GPU3D_Compute.cpp GPU3D_OpenGL_shaders.h OpenGLSupport.cpp) diff --git a/src/GPU3D.h b/src/GPU3D.h index 44d422a5..a20b0b33 100644 --- a/src/GPU3D.h +++ b/src/GPU3D.h @@ -169,6 +169,7 @@ extern std::unique_ptr CurrentRenderer; #ifdef OGLRENDERER_ENABLED #include "GPU3D_OpenGL.h" +#include "GPU3D_Compute.h" #endif #endif diff --git a/src/GPU3D_Compute.cpp b/src/GPU3D_Compute.cpp new file mode 100644 index 00000000..de8b3393 --- /dev/null +++ b/src/GPU3D_Compute.cpp @@ -0,0 +1,1451 @@ +/* + Copyright 2016-2022 melonDS team + + This file is part of melonDS. + + melonDS is free software: you can redistribute it and/or modify it under + the terms of the GNU General Public License as published by the Free + Software Foundation, either version 3 of the License, or (at your option) + any later version. + + melonDS is distributed in the hope that it will be useful, but WITHOUT ANY + WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS + FOR A PARTICULAR PURPOSE. See the GNU General Public License for more details. + + You should have received a copy of the GNU General Public License along + with melonDS. If not, see http://www.gnu.org/licenses/. +*/ + +#include "GPU3D_Compute.h" + +#include + +#define XXH_STATIC_LINKING_ONLY +#include "xxhash/xxhash.h" + +#include "OpenGLSupport.h" + +#include "GPU3D_Compute_shaders.h" + +namespace GPU3D +{ + +ComputeRenderer::ComputeRenderer() + : Renderer3D(true) +{} + +ComputeRenderer::~ComputeRenderer() +{} + + + +bool ComputeRenderer::CompileShader(GLuint& shader, const char* source, const std::initializer_list& defines) +{ + std::string shaderName; + std::string shaderSource; + shaderSource += "#version 430 core\n"; + for (const char* define : defines) + { + shaderSource += "#define "; + shaderSource += define; + shaderSource += '\n'; + shaderName += define; + shaderName += ','; + } + shaderSource += ComputeRendererShaders::Common; + shaderSource += source; + + return OpenGL::CompileComputeProgram(shader, shaderSource.c_str(), shaderName.c_str()); +} + +void blah(GLenum source,GLenum type,GLuint id,GLenum severity,GLsizei length,const GLchar *message,const void *userParam) +{ + printf("%s\n", message); +} + +bool ComputeRenderer::Init() +{ + glDebugMessageCallback(blah, NULL); + glEnable(GL_DEBUG_OUTPUT); + glGenBuffers(1, &YSpanSetupMemory); + glBindBuffer(GL_SHADER_STORAGE_BUFFER, YSpanSetupMemory); + glBufferData(GL_SHADER_STORAGE_BUFFER, sizeof(SpanSetupY)*MaxYSpanSetups, nullptr, GL_DYNAMIC_DRAW); + + glGenBuffers(1, &RenderPolygonMemory); + glBindBuffer(GL_SHADER_STORAGE_BUFFER, RenderPolygonMemory); + glBufferData(GL_SHADER_STORAGE_BUFFER, sizeof(RenderPolygon)*2048, nullptr, GL_DYNAMIC_DRAW); + + glGenBuffers(1, &TileMemory); + glBindBuffer(GL_SHADER_STORAGE_BUFFER, TileMemory); + glBufferData(GL_SHADER_STORAGE_BUFFER, sizeof(Tiles), nullptr, GL_DYNAMIC_DRAW); + + glGenBuffers(1, &XSpanSetupMemory); + glBindBuffer(GL_SHADER_STORAGE_BUFFER, XSpanSetupMemory); + glBufferData(GL_SHADER_STORAGE_BUFFER, sizeof(SpanSetupX)*MaxYSpanIndices, nullptr, GL_DYNAMIC_DRAW); + + glGenBuffers(1, &BinResultMemory); + glBindBuffer(GL_SHADER_STORAGE_BUFFER, BinResultMemory); + glBufferData(GL_SHADER_STORAGE_BUFFER, sizeof(BinResult), nullptr, GL_DYNAMIC_DRAW); + + glGenBuffers(1, &FinalTileMemory); + glBindBuffer(GL_SHADER_STORAGE_BUFFER, FinalTileMemory); + glBufferData(GL_SHADER_STORAGE_BUFFER, sizeof(FinalTiles), nullptr, GL_DYNAMIC_DRAW); + + glGenBuffers(1, &YSpanIndicesTextureMemory); + glBindBuffer(GL_TEXTURE_BUFFER, YSpanIndicesTextureMemory); + glBufferData(GL_TEXTURE_BUFFER, MaxYSpanIndices*2*4, nullptr, GL_DYNAMIC_DRAW); + + glGenTextures(1, &YSpanIndicesTexture); + glBindTexture(GL_TEXTURE_BUFFER, YSpanIndicesTexture); + glTexBuffer(GL_TEXTURE_BUFFER, GL_RGBA16UI, YSpanIndicesTextureMemory); + + CompileShader(ShaderInterpXSpans[0], ComputeRendererShaders::InterpSpans, {"InterpSpans", "ZBuffer"}); + CompileShader(ShaderInterpXSpans[1], ComputeRendererShaders::InterpSpans, {"InterpSpans", "WBuffer"}); + CompileShader(ShaderBinCombined, ComputeRendererShaders::BinCombined, {"BinCombined"}); + CompileShader(ShaderDepthBlend[0], ComputeRendererShaders::DepthBlend, {"DepthBlend", "ZBuffer"}); + CompileShader(ShaderDepthBlend[1], ComputeRendererShaders::DepthBlend, {"DepthBlend", "WBuffer"}); + CompileShader(ShaderRasteriseNoTexture[0], ComputeRendererShaders::Rasterise, {"Rasterise", "ZBuffer", "NoTexture"}); + CompileShader(ShaderRasteriseNoTexture[1], ComputeRendererShaders::Rasterise, {"Rasterise", "WBuffer", "NoTexture"}); + CompileShader(ShaderRasteriseNoTextureToon[0], ComputeRendererShaders::Rasterise, {"Rasterise", "ZBuffer", "NoTexture", "Toon"}); + CompileShader(ShaderRasteriseNoTextureToon[1], ComputeRendererShaders::Rasterise, {"Rasterise", "WBuffer", "NoTexture", "Toon"}); + CompileShader(ShaderRasteriseNoTextureHighlight[0], ComputeRendererShaders::Rasterise, {"Rasterise", "ZBuffer", "NoTexture", "Highlight"}); + CompileShader(ShaderRasteriseNoTextureHighlight[1], ComputeRendererShaders::Rasterise, {"Rasterise", "WBuffer", "NoTexture", "Highlight"}); + CompileShader(ShaderRasteriseUseTextureDecal[0], ComputeRendererShaders::Rasterise, {"Rasterise", "ZBuffer", "UseTexture", "Decal"}); + CompileShader(ShaderRasteriseUseTextureDecal[1], ComputeRendererShaders::Rasterise, {"Rasterise", "WBuffer", "UseTexture", "Decal"}); + CompileShader(ShaderRasteriseUseTextureModulate[0], ComputeRendererShaders::Rasterise, {"Rasterise", "ZBuffer", "UseTexture", "Modulate"}); + CompileShader(ShaderRasteriseUseTextureModulate[1], ComputeRendererShaders::Rasterise, {"Rasterise", "WBuffer", "UseTexture", "Modulate"}); + CompileShader(ShaderRasteriseUseTextureToon[0], ComputeRendererShaders::Rasterise, {"Rasterise", "ZBuffer", "UseTexture", "Toon"}); + CompileShader(ShaderRasteriseUseTextureToon[1], ComputeRendererShaders::Rasterise, {"Rasterise", "WBuffer", "UseTexture", "Toon"}); + CompileShader(ShaderRasteriseUseTextureHighlight[0], ComputeRendererShaders::Rasterise, {"Rasterise", "ZBuffer", "UseTexture", "Highlight"}); + CompileShader(ShaderRasteriseUseTextureHighlight[1], ComputeRendererShaders::Rasterise, {"Rasterise", "WBuffer", "UseTexture", "Highlight"}); + CompileShader(ShaderRasteriseShadowMask[0], ComputeRendererShaders::Rasterise, {"Rasterise", "ZBuffer", "ShadowMask"}); + CompileShader(ShaderRasteriseShadowMask[1], ComputeRendererShaders::Rasterise, {"Rasterise", "WBuffer", "ShadowMask"}); + CompileShader(ShaderClearCoarseBinMask, ComputeRendererShaders::ClearCoarseBinMask, {"ClearCoarseBinMask"}); + CompileShader(ShaderClearIndirectWorkCount, ComputeRendererShaders::ClearIndirectWorkCount, {"ClearIndirectWorkCount"}); + CompileShader(ShaderCalculateWorkListOffset, ComputeRendererShaders::CalcOffsets, {"CalculateWorkOffsets"}); + CompileShader(ShaderSortWork, ComputeRendererShaders::SortWork, {"SortWork"}); + CompileShader(ShaderFinalPass[0], ComputeRendererShaders::FinalPass, {"FinalPass"}); + CompileShader(ShaderFinalPass[1], ComputeRendererShaders::FinalPass, {"FinalPass", "EdgeMarking"}); + CompileShader(ShaderFinalPass[2], ComputeRendererShaders::FinalPass, {"FinalPass", "Fog"}); + CompileShader(ShaderFinalPass[3], ComputeRendererShaders::FinalPass, {"FinalPass", "EdgeMarking", "Fog"}); + CompileShader(ShaderFinalPass[4], ComputeRendererShaders::FinalPass, {"FinalPass", "AntiAliasing"}); + CompileShader(ShaderFinalPass[5], ComputeRendererShaders::FinalPass, {"FinalPass", "AntiAliasing", "EdgeMarking"}); + CompileShader(ShaderFinalPass[6], ComputeRendererShaders::FinalPass, {"FinalPass", "AntiAliasing", "Fog"}); + CompileShader(ShaderFinalPass[7], ComputeRendererShaders::FinalPass, {"FinalPass", "AntiAliasing", "EdgeMarking", "Fog"}); + + return true; +} + +void ComputeRenderer::DeInit() +{ + +} + +void ComputeRenderer::Reset() +{ + /*for (u32 i = 0; i < 8; i++) + { + for (u32 j = 0; j < 8; j++) + { + for (u32 k = 0; k < TexArrays[i][j].size(); k++) + Gfx::TextureHeap->Free(TexArrays[i][j][k].Memory); + TexArrays[i][j].clear(); + FreeTextures[i][j].clear(); + } + }*/ + TexCache.clear(); + + FreeImageDescriptorsCount = TexCacheMaxImages; + for (int i = 0; i < TexCacheMaxImages; i++) + { + FreeImageDescriptors[i] = i; + } +} + +void ComputeRenderer::SetRenderSettings(GPU::RenderSettings& settings) +{ + +} + +void ComputeRenderer::VCount144() +{ + +} + +void ComputeRenderer::SetupAttrs(SpanSetupY* span, Polygon* poly, int from, int to) +{ + span->Z0 = poly->FinalZ[from]; + span->W0 = poly->FinalW[from]; + span->Z1 = poly->FinalZ[to]; + span->W1 = poly->FinalW[to]; + span->ColorR0 = poly->Vertices[from]->FinalColor[0]; + span->ColorG0 = poly->Vertices[from]->FinalColor[1]; + span->ColorB0 = poly->Vertices[from]->FinalColor[2]; + span->ColorR1 = poly->Vertices[to]->FinalColor[0]; + span->ColorG1 = poly->Vertices[to]->FinalColor[1]; + span->ColorB1 = poly->Vertices[to]->FinalColor[2]; + span->TexcoordU0 = poly->Vertices[from]->TexCoords[0]; + span->TexcoordV0 = poly->Vertices[from]->TexCoords[1]; + span->TexcoordU1 = poly->Vertices[to]->TexCoords[0]; + span->TexcoordV1 = poly->Vertices[to]->TexCoords[1]; +} + +void ComputeRenderer::SetupYSpanDummy(SpanSetupY* span, Polygon* poly, int vertex, int side) +{ + s32 x0 = poly->Vertices[vertex]->FinalPosition[0]; + if (side) + { + span->DxInitial = -0x40000; + x0--; + } + else + { + span->DxInitial = 0; + } + + span->X0 = span->X1 = x0; + span->XMin = x0; + span->XMax = x0; + span->Y0 = span->Y1 = poly->Vertices[vertex]->FinalPosition[1]; + + span->Increment = 0; + + span->I0 = span->I1 = span->IRecip = 0; + span->Linear = true; + + span->XCovIncr = 0; + + span->IsDummy = true; + + SetupAttrs(span, poly, vertex, vertex); +} + +void ComputeRenderer::SetupYSpan(int polynum, SpanSetupY* span, Polygon* poly, int from, int to, u32 y, int side) +{ + span->X0 = poly->Vertices[from]->FinalPosition[0]; + span->X1 = poly->Vertices[to]->FinalPosition[0]; + span->Y0 = poly->Vertices[from]->FinalPosition[1]; + span->Y1 = poly->Vertices[to]->FinalPosition[1]; + + SetupAttrs(span, poly, from, to); + + bool negative = false; + if (span->X1 > span->X0) + { + span->XMin = span->X0; + span->XMax = span->X1-1; + } + else if (span->X1 < span->X0) + { + span->XMin = span->X1; + span->XMax = span->X0-1; + negative = true; + } + else + { + span->XMin = span->X0; + if (side) span->XMin--; + span->XMax = span->XMin; + } + + span->IsDummy = false; + + s32 xlen = span->XMax+1 - span->XMin; + s32 ylen = span->Y1 - span->Y0; + + // slope increment has a 18-bit fractional part + // note: for some reason, x/y isn't calculated directly, + // instead, 1/y is calculated and then multiplied by x + // TODO: this is still not perfect (see for example x=169 y=33) + if (ylen == 0) + { + span->Increment = 0; + } + else if (ylen == xlen) + { + span->Increment = 0x40000; + } + else + { + s32 yrecip = (1<<18) / ylen; + span->Increment = (span->X1-span->X0) * yrecip; + if (span->Increment < 0) span->Increment = -span->Increment; + } + + bool xMajor = (span->Increment > 0x40000); + + if (side) + { + // right + + if (xMajor) + span->DxInitial = negative ? (0x20000 + 0x40000) : (span->Increment - 0x20000); + else if (span->Increment != 0) + span->DxInitial = negative ? 0x40000 : 0; + else + span->DxInitial = -0x40000; + } + else + { + // left + + if (xMajor) + span->DxInitial = negative ? ((span->Increment - 0x20000) + 0x40000) : 0x20000; + else if (span->Increment != 0) + span->DxInitial = negative ? 0x40000 : 0; + else + span->DxInitial = 0; + } + + if (xMajor) + { + if (side) + { + span->I0 = span->X0 - 1; + span->I1 = span->X1 - 1; + } + else + { + span->I0 = span->X0; + span->I1 = span->X1; + } + + // used for calculating AA coverage + span->XCovIncr = (ylen << 10) / xlen; + } + else + { + span->I0 = span->Y0; + span->I1 = span->Y1; + } + + //if (span->I1 < span->I0) + // std::swap(span->I0, span->I1); + + if (span->I0 != span->I1) + span->IRecip = (1<<30) / (span->I1 - span->I0); + else + span->IRecip = 0; + + span->Linear = (span->W0 == span->W1) && !(span->W0 & 0x7E) && !(span->W1 & 0x7E); + + if ((span->W0 & 0x1) && !(span->W1 & 0x1)) + { + span->W0n = (span->W0 - 1) >> 1; + span->W0d = (span->W0 + 1) >> 1; + span->W1d = span->W1 >> 1; + } + else + { + span->W0n = span->W0 >> 1; + span->W0d = span->W0 >> 1; + span->W1d = span->W1 >> 1; + } +} + +inline u32 TextureWidth(u32 texparam) +{ + return 8 << ((texparam >> 20) & 0x7); +} + +inline u32 TextureHeight(u32 texparam) +{ + return 8 << ((texparam >> 23) & 0x7); +} + +inline u16 ColorAvg(u16 color0, u16 color1) +{ + u32 r0 = color0 & 0x001F; + u32 g0 = color0 & 0x03E0; + u32 b0 = color0 & 0x7C00; + u32 r1 = color1 & 0x001F; + u32 g1 = color1 & 0x03E0; + u32 b1 = color1 & 0x7C00; + + u32 r = (r0 + r1) >> 1; + u32 g = ((g0 + g1) >> 1) & 0x03E0; + u32 b = ((b0 + b1) >> 1) & 0x7C00; + + return r | g | b; +} + +inline u16 Color5of3(u16 color0, u16 color1) +{ + u32 r0 = color0 & 0x001F; + u32 g0 = color0 & 0x03E0; + u32 b0 = color0 & 0x7C00; + u32 r1 = color1 & 0x001F; + u32 g1 = color1 & 0x03E0; + u32 b1 = color1 & 0x7C00; + + u32 r = (r0*5 + r1*3) >> 3; + u32 g = ((g0*5 + g1*3) >> 3) & 0x03E0; + u32 b = ((b0*5 + b1*3) >> 3) & 0x7C00; + + return r | g | b; +} + +inline u16 Color3of5(u16 color0, u16 color1) +{ + u32 r0 = color0 & 0x001F; + u32 g0 = color0 & 0x03E0; + u32 b0 = color0 & 0x7C00; + u32 r1 = color1 & 0x001F; + u32 g1 = color1 & 0x03E0; + u32 b1 = color1 & 0x7C00; + + u32 r = (r0*3 + r1*5) >> 3; + u32 g = ((g0*3 + g1*5) >> 3) & 0x03E0; + u32 b = ((b0*3 + b1*5) >> 3) & 0x7C00; + + return r | g | b; +} + +/* +inline void RGB5ToRGB6(uint8x16_t lo, uint8x16_t hi, uint8x16_t& red, uint8x16_t& green, uint8x16_t& blue) +{ + red = vandq_u8(vshlq_n_u8(lo, 1), vdupq_n_u8(0x3E)); + green = vbslq_u8(vdupq_n_u8(0xCE), vshrq_n_u8(lo, 4), vshlq_n_u8(hi, 4)); + blue = vandq_u8(vshrq_n_u8(hi, 1), vdupq_n_u8(0x3E)); + red = vandq_u8(vtstq_u8(red, red), vaddq_u8(red, vdupq_n_u8(1))); + green = vandq_u8(vtstq_u8(green, green), vaddq_u8(green, vdupq_n_u8(1))); + blue = vandq_u8(vtstq_u8(blue, blue), vaddq_u8(blue, vdupq_n_u8(1))); +} + +inline void RGB5ToRGB6(uint8x8_t lo, uint8x8_t hi, uint8x8_t& red, uint8x8_t& green, uint8x8_t& blue) +{ + red = vand_u8(vshl_n_u8(lo, 1), vdup_n_u8(0x3E)); + green = vbsl_u8(vdup_n_u8(0xCE), vshr_n_u8(lo, 4), vshl_n_u8(hi, 4)); + blue = vand_u8(vshr_n_u8(hi, 1), vdup_n_u8(0x3E)); + + red = vand_u8(vtst_u8(red, red), vadd_u8(red, vdup_n_u8(1))); + green = vand_u8(vtst_u8(green, green), vadd_u8(green, vdup_n_u8(1))); + blue = vand_u8(vtst_u8(blue, blue), vadd_u8(blue, vdup_n_u8(1))); +}*/ + +inline u32 ConvertRGB5ToRGB8(u16 val) +{ + return (((u32)val & 0x1F) << 3) + | (((u32)val & 0x3E0) << 6) + | (((u32)val & 0x7C00) << 9); +} +inline u32 ConvertRGB5ToBGR8(u16 val) +{ + return (((u32)val & 0x1F) << 9) + | (((u32)val & 0x3E0) << 6) + | (((u32)val & 0x7C00) << 3); +} +inline u32 ConvertRGB5ToRGB6(u16 val) +{ + u8 r = (val & 0x1F) << 1; + u8 g = (val & 0x3E0) >> 4; + u8 b = (val & 0x7C00) >> 9; + if (r) r++; + if (g) g++; + if (b) b++; + return (u32)r | ((u32)g << 8) | ((u32)b << 16); +} + +enum +{ + outputFmt_RGB6A5, + outputFmt_RGBA8, + outputFmt_BGRA8 +}; + +template +void ConvertCompressedTexture(u32 width, u32 height, u32* output, u8* texData, u8* texAuxData, u16* palData) +{ + // we process a whole block at the time + for (int y = 0; y < height / 4; y++) + { + for (int x = 0; x < width / 4; x++) + { + u32 data = ((u32*)texData)[x + y * (width / 4)]; + u16 auxData = ((u16*)texAuxData)[x + y * (width / 4)]; + + u32 paletteOffset = auxData & 0x3FFF; + u16 color0 = palData[paletteOffset*2] | 0x8000; + u16 color1 = palData[paletteOffset*2+1] | 0x8000; + u16 color2, color3; + + switch ((auxData >> 14) & 0x3) + { + case 0: + color2 = palData[paletteOffset*2+2] | 0x8000; + color3 = 0; + break; + case 1: + { + u32 r0 = color0 & 0x001F; + u32 g0 = color0 & 0x03E0; + u32 b0 = color0 & 0x7C00; + u32 r1 = color1 & 0x001F; + u32 g1 = color1 & 0x03E0; + u32 b1 = color1 & 0x7C00; + + u32 r = (r0 + r1) >> 1; + u32 g = ((g0 + g1) >> 1) & 0x03E0; + u32 b = ((b0 + b1) >> 1) & 0x7C00; + color2 = r | g | b | 0x8000; + } + color3 = 0; + break; + case 2: + color2 = palData[paletteOffset*2+2] | 0x8000; + color3 = palData[paletteOffset*2+3] | 0x8000; + break; + case 3: + { + u32 r0 = color0 & 0x001F; + u32 g0 = color0 & 0x03E0; + u32 b0 = color0 & 0x7C00; + u32 r1 = color1 & 0x001F; + u32 g1 = color1 & 0x03E0; + u32 b1 = color1 & 0x7C00; + + u32 r = (r0*5 + r1*3) >> 3; + u32 g = ((g0*5 + g1*3) >> 3) & 0x03E0; + u32 b = ((b0*5 + b1*3) >> 3) & 0x7C00; + + color2 = r | g | b | 0x8000; + } + { + u32 r0 = color0 & 0x001F; + u32 g0 = color0 & 0x03E0; + u32 b0 = color0 & 0x7C00; + u32 r1 = color1 & 0x001F; + u32 g1 = color1 & 0x03E0; + u32 b1 = color1 & 0x7C00; + + u32 r = (r0*3 + r1*5) >> 3; + u32 g = ((g0*3 + g1*5) >> 3) & 0x03E0; + u32 b = ((b0*3 + b1*5) >> 3) & 0x7C00; + + color3 = r | g | b | 0x8000; + } + break; + } + + // in 2020 our default data types are big enough to be used as lookup tables... + u64 packed = color0 | ((u64)color1 << 16) | ((u64)color2 << 32) | ((u64)color3 << 48); + + for (int j = 0; j < 4; j++) + { + for (int i = 0; i < 4; i++) + { + u16 color = (packed >> 16 * (data >> 2 * (i + j * 4))) & 0xFFFF; + u32 res; + switch (outputFmt) + { + case outputFmt_RGB6A5: res = ConvertRGB5ToRGB6(color) + | ((color & 0x8000) ? 0x1F000000 : 0); break; + case outputFmt_RGBA8: res = ConvertRGB5ToRGB8(color) + | ((color & 0x8000) ? 0xFF000000 : 0); break; + case outputFmt_BGRA8: res = ConvertRGB5ToBGR8(color) + | ((color & 0x8000) ? 0xFF000000 : 0); break; + } + output[x * 4 + i + (y * 4 + j) * width] = res; + } + } + } + } +} + +template +void ConvertAXIYTexture(u32 width, u32 height, u32* output, u8* texData, u16* palData) +{ + /*for (int y = 0; y < height; y++) + { + for (int x = 0; x < width; x++) + { + u8 val = texData[x + y * width]; + + u32 idx = val & ((1 << Y) - 1); + + u16 color = palData[idx]; + u32 alpha = (val >> Y) & ((1 << X) - 1); + if (X != 5) + alpha = alpha * 4 + alpha / 2; + + u32 res; + switch (outputFmt) + { + case outputFmt_RGB6A5: res = ConvertRGB5ToRGB6(color) | alpha << 24; break; + // make sure full alpha == 255 + case outputFmt_RGBA8: res = ConvertRGB5ToRGB8(color) | (alpha << 27 | (alpha & 0x1C) << 22); break; + case outputFmt_BGRA8: res = ConvertRGB5ToBGR8(color) | (alpha << 27 | (alpha & 0x1C) << 22); break; + } + output[x + y * width] = res; + } + }*/ +} + +void Convert16ColorsTexture(u32 width, u32 height, u32* output, u8* texData, u16* palData, bool color0Transparent) +{ + /*uint8x16x2_t palette = vld2q_u8((u8*)palData); + + uint8x16_t paletteR, paletteG, paletteB; + RGB5ToRGB6(palette.val[0], palette.val[1], paletteR, paletteG, paletteB); + + uint8x16_t firstEntryAlpha = vdupq_n_u8(color0Transparent ? 0 : 0x1F); + + for (int i = 0; i < width*height/2; i += 16) + { + uint8x16_t packedIndices = vld1q_u8(&texData[i]); + + // unpack indices + uint8x16_t oddIndices = vandq_u8(packedIndices, vdupq_n_u8(0xF)); + uint8x16_t evenIndices = vshrq_n_u8(packedIndices, 4); + + uint8x16_t indices0 = vzip1q_u8(oddIndices, evenIndices); + uint8x16_t indices1 = vzip2q_u8(oddIndices, evenIndices); + + // palettise + uint8x16x4_t finalPixels0, finalPixels1; + finalPixels0.val[0] = vqtbl1q_u8(paletteR, indices0); + finalPixels0.val[1] = vqtbl1q_u8(paletteG, indices0); + finalPixels0.val[2] = vqtbl1q_u8(paletteB, indices0); + finalPixels0.val[3] = vbslq_u8(vceqzq_u8(indices0), firstEntryAlpha, vdupq_n_u8(0x1F)); + finalPixels1.val[0] = vqtbl1q_u8(paletteR, indices1); + finalPixels1.val[1] = vqtbl1q_u8(paletteG, indices1); + finalPixels1.val[2] = vqtbl1q_u8(paletteB, indices1); + finalPixels1.val[3] = vbslq_u8(vceqzq_u8(indices1), firstEntryAlpha, vdupq_n_u8(0x1F)); + + vst4q_u8((u8*)&output[i*2], finalPixels0); + vst4q_u8((u8*)&output[i*2+16], finalPixels1); + }*/ +} + +template +void ConvertNColorsTexture(u32 width, u32 height, u32* output, u8* texData, u16* palData, bool color0Transparent) +{ + for (int y = 0; y < height; y++) + { + for (int x = 0; x < width / (8 / colorBits); x++) + { + u8 val = texData[x + y * (width / (8 / colorBits))]; + + for (int i = 0; i < 8 / colorBits; i++) + { + u32 index = (val >> (i * colorBits)) & ((1 << colorBits) - 1); + u16 color = palData[index]; + + bool transparent = color0Transparent && index == 0; + u32 res; + switch (outputFmt) + { + case outputFmt_RGB6A5: res = ConvertRGB5ToRGB6(color) + | (transparent ? 0 : 0x1F000000); break; + case outputFmt_RGBA8: res = ConvertRGB5ToRGB8(color) + | (transparent ? 0 : 0xFF000000); break; + case outputFmt_BGRA8: res = ConvertRGB5ToBGR8(color) + | (transparent ? 0 : 0xFF000000); break; + } + output[x * (8 / colorBits) + y * width + i] = res; + } + } + } +} + +ComputeRenderer::TexCacheEntry& ComputeRenderer::GetTexture(u32 texParam, u32 palBase) +{ + // remove sampling and texcoord gen params + texParam &= ~0xC00F0000; + + u32 fmt = (texParam >> 26) & 0x7; + u64 key = texParam; + if (fmt != 7) + { + key |= (u64)palBase << 32; + if (fmt == 5) + key &= ~((u64)1 << 29); + } + //printf("%" PRIx64 " %" PRIx32 " %" PRIx32 "\n", key, texParam, palBase); + + assert(fmt != 0 && "no texture is not a texture format!"); + + auto it = TexCache.find(key); + + if (it != TexCache.end()) + return it->second; + + u32 widthLog2 = (texParam >> 20) & 0x7; + u32 heightLog2 = (texParam >> 23) & 0x7; + u32 width = 8 << widthLog2; + u32 height = 8 << heightLog2; + + u32 addr = (texParam & 0xFFFF) * 8; + + TexCacheEntry entry = {0}; + + entry.TextureRAMStart[0] = addr; + entry.WidthLog2 = widthLog2; + entry.HeightLog2 = heightLog2; + + // apparently a new texture + if (fmt == 7) + { + entry.TextureRAMSize[0] = width*height*2; + + /*for (u32 i = 0; i < width*height; i += 16) + { + uint8x16x2_t pixels = vld2q_u8(&GPU::VRAMFlat_Texture[addr + i * 2]); + + uint8x16_t red, green, blue; + RGB5ToRGB6(pixels.val[0], pixels.val[1], red, green, blue); + uint8x16_t alpha = vbslq_u8(vtstq_u8(pixels.val[1], vdupq_n_u8(0x80)), vdupq_n_u8(0x1F), vdupq_n_u8(0)); + + vst4q_u8((u8*)&TextureDecodingBuffer[i], + { + red, + green, + blue, + alpha + }); + }*/ + } + else if (fmt == 5) + { + u8* texData = &GPU::VRAMFlat_Texture[addr]; + u32 slot1addr = 0x20000 + ((addr & 0x1FFFC) >> 1); + if (addr >= 0x40000) + slot1addr += 0x10000; + u8* texAuxData = &GPU::VRAMFlat_Texture[slot1addr]; + + u16* palData = (u16*)(GPU::VRAMFlat_TexPal + palBase*16); + + entry.TextureRAMSize[0] = width*height/16*4; + entry.TextureRAMStart[1] = slot1addr; + entry.TextureRAMSize[1] = width*height/16*2; + entry.TexPalStart = palBase*16; + entry.TexPalSize = 0x10000; + + ConvertCompressedTexture(width, height, TextureDecodingBuffer, texData, texAuxData, palData); + } + else + { + u32 texSize, palAddr = palBase*16, numPalEntries; + switch (fmt) + { + case 1: texSize = width*height; numPalEntries = 32; break; + case 6: texSize = width*height; numPalEntries = 8; break; + case 2: texSize = width*height/4; numPalEntries = 4; palAddr >>= 1; break; + case 3: texSize = width*height/2; numPalEntries = 16; break; + case 4: texSize = width*height; numPalEntries = 256; break; + } + + palAddr &= 0x1FFFF; + + /*printf("creating texture | fmt: %d | %dx%d | %08x | %08x\n", fmt, width, height, addr, palAddr); + svcSleepThread(1000*1000);*/ + + entry.TextureRAMSize[0] = texSize; + entry.TexPalStart = palAddr; + entry.TexPalSize = numPalEntries*2; + + u8* texData = &GPU::VRAMFlat_Texture[addr]; + u16* palData = (u16*)(GPU::VRAMFlat_TexPal + palAddr); + + //assert(entry.TexPalStart+entry.TexPalSize <= 128*1024*1024); + + bool color0Transparent = texParam & (1 << 29); + + switch (fmt) + { + case 1: ConvertAXIYTexture(width, height, TextureDecodingBuffer, texData, palData); break; + case 6: ConvertAXIYTexture(width, height, TextureDecodingBuffer, texData, palData); break; + case 2: ConvertNColorsTexture(width, height, TextureDecodingBuffer, texData, palData, color0Transparent); break; + case 3: Convert16ColorsTexture(width, height, TextureDecodingBuffer, texData, palData, color0Transparent); break; + case 4: ConvertNColorsTexture(width, height, TextureDecodingBuffer, texData, palData, color0Transparent); break; + } + } + + for (int i = 0; i < 2; i++) + { + if (entry.TextureRAMSize[i]) + entry.TextureHash[i] = XXH3_64bits(&GPU::VRAMFlat_Texture[entry.TextureRAMStart[i]], entry.TextureRAMSize[i]); + } + if (entry.TexPalSize) + entry.TexPalHash = XXH3_64bits(&GPU::VRAMFlat_TexPal[entry.TexPalStart], entry.TexPalSize); + + auto& texArrays = TexArrays[widthLog2][heightLog2]; + auto& freeTextures = FreeTextures[widthLog2][heightLog2]; + + /*if (freeTextures.size() == 0) + { + texArrays.resize(texArrays.size()+1); + TexArray& array = texArrays[texArrays.size()-1]; + + u32 layers = std::min((8*1024*1024) / (width*height*4), 64); + + // allocate new array texture + dk::ImageLayout imageLayout; + dk::ImageLayoutMaker{Gfx::Device} + .setType(DkImageType_2DArray) + .setFormat(DkImageFormat_RGBA8_Uint) + .setDimensions(width, height, layers) + .initialize(imageLayout); + + assert(FreeImageDescriptorsCount > 0); + array.ImageDescriptor = FreeImageDescriptors[--FreeImageDescriptorsCount]; + + array.Memory = Gfx::TextureHeap->Alloc(imageLayout.getSize(), imageLayout.getAlignment()); + array.Image.initialize(imageLayout, Gfx::TextureHeap->MemBlock, array.Memory.Offset); + + dk::ImageDescriptor descriptor; + descriptor.initialize(array.Image); + DkGpuAddr descriptors = Gfx::DataHeap->GpuAddr(ImageDescriptors); + EmuCmdBuf.pushData(descriptors + (descriptorOffset_TexcacheStart + array.ImageDescriptor) * sizeof(DkImageDescriptor), + &descriptor, + sizeof(DkImageDescriptor)); + + //printf("allocating new layer set for %d %d %d %d\n", width, height, texArrays.size()-1, array.ImageDescriptor); + + for (u16 i = 0; i < layers; i++) + { + freeTextures.push_back(TexArrayEntry{(u16)(texArrays.size()-1), i}); + } + }*/ + + TexArrayEntry storagePlace = freeTextures[freeTextures.size()-1]; + freeTextures.pop_back(); + + TexArray& array = texArrays[storagePlace.TexArrayIdx]; + //printf("using storage place %d %d | %d %d (%d)\n", width, height, storagePlace.TexArrayIdx, storagePlace.LayerIdx, array.ImageDescriptor); + + /*UploadBuf.UploadAndCopyTexture(Gfx::EmuCmdBuf, array.Image, + (u8*)TextureDecodingBuffer, + 0, 0, width, height, + width*4, + storagePlace.LayerIdx);*/ + + entry.DescriptorIdx = array.ImageDescriptor; + entry.Texture = storagePlace; + + return TexCache.emplace(std::make_pair(key, entry)).first->second; +} + +struct Variant +{ + s16 Texture, Sampler; + u16 Width, Height; + u8 BlendMode; + + bool operator==(const Variant& other) + { + return Texture == other.Texture && Sampler == other.Sampler && BlendMode == other.BlendMode; + } +}; + +/* + Antialiasing + W-Buffer + Mit Textur + 0 + 1, 3 + 2 + Ohne Textur + 2 + 0, 1, 3 + + => 20 Shader + 1x Shadow Mask +*/ + +void ComputeRenderer::RenderFrame() +{ + //printf("render frame\n"); + auto textureDirty = GPU::VRAMDirty_Texture.DeriveState(GPU::VRAMMap_Texture); + auto texPalDirty = GPU::VRAMDirty_TexPal.DeriveState(GPU::VRAMMap_TexPal); + + bool textureChanged = GPU::MakeVRAMFlat_TextureCoherent(textureDirty); + bool texPalChanged = GPU::MakeVRAMFlat_TexPalCoherent(texPalDirty); + + if (textureChanged || texPalChanged) + { + //printf("check invalidation %d\n", TexCache.size()); + for (auto it = TexCache.begin(); it != TexCache.end();) + { + TexCacheEntry& entry = it->second; + if (textureChanged) + { + for (u32 i = 0; i < 2; i++) + { + u32 startBit = entry.TextureRAMStart[i] / GPU::VRAMDirtyGranularity; + u32 bitsCount = ((entry.TextureRAMStart[i] + entry.TextureRAMSize[i] + GPU::VRAMDirtyGranularity - 1) / GPU::VRAMDirtyGranularity) - startBit; + + u32 startEntry = startBit >> 6; + u64 entriesCount = ((startBit + bitsCount + 0x3F) >> 6) - startEntry; + for (u32 j = startEntry; j < startEntry + entriesCount; j++) + { + if (GetRangedBitMask(j, startBit, bitsCount) & textureDirty.Data[j]) + { + u64 newTexHash = XXH3_64bits(&GPU::VRAMFlat_Texture[entry.TextureRAMStart[i]], entry.TextureRAMSize[i]); + + if (newTexHash != entry.TextureHash[i]) + goto invalidate; + } + } + } + } + + if (texPalChanged && entry.TexPalSize > 0) + { + u32 startBit = entry.TexPalStart / GPU::VRAMDirtyGranularity; + u32 bitsCount = ((entry.TexPalStart + entry.TexPalSize + GPU::VRAMDirtyGranularity - 1) / GPU::VRAMDirtyGranularity) - startBit; + + u32 startEntry = startBit >> 6; + u64 entriesCount = ((startBit + bitsCount + 0x3F) >> 6) - startEntry; + for (u32 j = startEntry; j < startEntry + entriesCount; j++) + { + if (GetRangedBitMask(j, startBit, bitsCount) & texPalDirty.Data[j]) + { + u64 newPalHash = XXH3_64bits(&GPU::VRAMFlat_TexPal[entry.TexPalStart], entry.TexPalSize); + if (newPalHash != entry.TexPalHash) + goto invalidate; + } + } + } + + it++; + continue; + invalidate: + FreeTextures[entry.WidthLog2][entry.HeightLog2].push_back(entry.Texture); + + //printf("invalidating texture %d\n", entry.ImageDescriptor); + + it = TexCache.erase(it); + } + } + else if (RenderFrameIdentical) + { + return; + } + + int numYSpans = 0; + int numSetupIndices = 0; + + u32 numVariants = 0, prevVariant, prevTexLayer; + Variant variants[MaxVariants]; + + int foundviatexcache = 0, foundviaprev = 0, numslow = 0; + + bool enableTextureMaps = RenderDispCnt & (1<<0); + + for (int i = 0; i < RenderNumPolygons; i++) + { + Polygon* polygon = RenderPolygonRAM[i]; + + u32 nverts = polygon->NumVertices; + u32 vtop = polygon->VTop, vbot = polygon->VBottom; + s32 ytop = polygon->YTop, ybot = polygon->YBottom; + + u32 curVL = vtop, curVR = vtop; + u32 nextVL, nextVR; + + RenderPolygons[i].FirstXSpan = numSetupIndices; + RenderPolygons[i].YTop = ytop; + RenderPolygons[i].YBot = ybot; + RenderPolygons[i].Attr = polygon->Attr; + + bool foundVariant = false; + if (i > 0) + { + Polygon* prevPolygon = RenderPolygonRAM[i - 1]; + foundVariant = prevPolygon->TexParam == polygon->TexParam + && prevPolygon->TexPalette == polygon->TexPalette + && (prevPolygon->Attr & 0x30) == (polygon->Attr & 0x30) + && prevPolygon->IsShadowMask == polygon->IsShadowMask; + if (foundVariant) + foundviaprev++; + } + + if (!foundVariant) + { + Variant variant; + variant.BlendMode = polygon->IsShadowMask ? 4 : ((polygon->Attr >> 4) & 0x3); + variant.Texture = -1; + variant.Sampler = -1; + TexCacheEntry* texcacheEntry = nullptr; + if (enableTextureMaps && (polygon->TexParam >> 26) & 0x7) + { + texcacheEntry = &GetTexture(polygon->TexParam, polygon->TexPalette); + bool wrapS = (polygon->TexParam >> 16) & 1; + bool wrapT = (polygon->TexParam >> 17) & 1; + bool mirrorS = (polygon->TexParam >> 18) & 1; + bool mirrorT = (polygon->TexParam >> 19) & 1; + variant.Sampler = (wrapS ? (mirrorS ? 2 : 1) : 0) + (wrapT ? (mirrorT ? 2 : 1) : 0) * 3; + variant.Texture = texcacheEntry->DescriptorIdx; + prevTexLayer = texcacheEntry->Texture.LayerIdx; + if (texcacheEntry->LastVariant < numVariants && variants[texcacheEntry->LastVariant] == variant) + { + foundVariant = true; + prevVariant = texcacheEntry->LastVariant; + foundviatexcache++; + } + } + + if (!foundVariant) + { + numslow++; + for (int j = numVariants - 1; j >= 0; j--) + { + if (variants[j] == variant) + { + foundVariant = true; + prevVariant = j; + goto foundVariant; + } + } + + prevVariant = numVariants; + variants[numVariants] = variant; + variants[numVariants].Width = TextureWidth(polygon->TexParam); + variants[numVariants].Height = TextureHeight(polygon->TexParam); + numVariants++; + assert(numVariants <= MaxVariants); + foundVariant:; + + if (texcacheEntry) + texcacheEntry->LastVariant = prevVariant; + } + } + RenderPolygons[i].Variant = prevVariant; + RenderPolygons[i].TextureLayer = (float)prevTexLayer; + + if (polygon->FacingView) + { + nextVL = curVL + 1; + if (nextVL >= nverts) nextVL = 0; + nextVR = curVR - 1; + if ((s32)nextVR < 0) nextVR = nverts - 1; + } + else + { + nextVL = curVL - 1; + if ((s32)nextVL < 0) nextVL = nverts - 1; + nextVR = curVR + 1; + if (nextVR >= nverts) nextVR = 0; + } + + s32 minX = polygon->Vertices[vtop]->FinalPosition[0]; + s32 minXY = polygon->Vertices[vtop]->FinalPosition[1]; + s32 maxX = polygon->Vertices[vtop]->FinalPosition[0]; + s32 maxXY = polygon->Vertices[vtop]->FinalPosition[1]; + + if (ybot == ytop) + { + vtop = 0; vbot = 0; + + RenderPolygons[i].YBot++; + + int j = 1; + if (polygon->Vertices[j]->FinalPosition[0] < polygon->Vertices[vtop]->FinalPosition[0]) vtop = j; + if (polygon->Vertices[j]->FinalPosition[0] > polygon->Vertices[vbot]->FinalPosition[0]) vbot = j; + + j = nverts - 1; + if (polygon->Vertices[j]->FinalPosition[0] < polygon->Vertices[vtop]->FinalPosition[0]) vtop = j; + if (polygon->Vertices[j]->FinalPosition[0] > polygon->Vertices[vbot]->FinalPosition[0]) vbot = j; + + assert(numYSpans < MaxYSpanSetups); + u32 curSpanL = numYSpans; + SetupYSpanDummy(&YSpanSetups[numYSpans++], polygon, vtop, 0); + assert(numYSpans < MaxYSpanSetups); + u32 curSpanR = numYSpans; + SetupYSpanDummy(&YSpanSetups[numYSpans++], polygon, vbot, 1); + + minX = YSpanSetups[curSpanL].X0; + minXY = YSpanSetups[curSpanL].Y0; + maxX = YSpanSetups[curSpanR].X0; + maxXY = YSpanSetups[curSpanR].Y0; + if (maxX < minX) + { + std::swap(minX, maxX); + std::swap(minXY, maxXY); + } + + assert(numSetupIndices < MaxYSpanIndices); + YSpanIndices[numSetupIndices].PolyIdx = i; + YSpanIndices[numSetupIndices].SpanIdxL = curSpanL; + YSpanIndices[numSetupIndices].SpanIdxR = curSpanR; + YSpanIndices[numSetupIndices].Y = ytop; + numSetupIndices++; + } + else + { + u32 curSpanL = numYSpans; + assert(numYSpans < MaxYSpanSetups); + SetupYSpan(i, &YSpanSetups[numYSpans++], polygon, curVL, nextVL, ytop, 0); + u32 curSpanR = numYSpans; + assert(numYSpans < MaxYSpanSetups); + SetupYSpan(i, &YSpanSetups[numYSpans++], polygon, curVR, nextVR, ytop, 1); + + for (u32 y = ytop; y < ybot; y++) + { + if (y >= polygon->Vertices[nextVL]->FinalPosition[1] && curVL != polygon->VBottom) + { + while (y >= polygon->Vertices[nextVL]->FinalPosition[1] && curVL != polygon->VBottom) + { + curVL = nextVL; + if (polygon->FacingView) + { + nextVL = curVL + 1; + if (nextVL >= nverts) + nextVL = 0; + } + else + { + nextVL = curVL - 1; + if ((s32)nextVL < 0) + nextVL = nverts - 1; + } + } + + if (polygon->Vertices[curVL]->FinalPosition[0] < minX) + { + minX = polygon->Vertices[curVL]->FinalPosition[0]; + minXY = polygon->Vertices[curVL]->FinalPosition[1]; + } + if (polygon->Vertices[curVL]->FinalPosition[0] > maxX) + { + maxX = polygon->Vertices[curVL]->FinalPosition[0]; + maxXY = polygon->Vertices[curVL]->FinalPosition[1]; + } + + assert(numYSpans < MaxYSpanSetups); + curSpanL = numYSpans; + SetupYSpan(i,&YSpanSetups[numYSpans++], polygon, curVL, nextVL, y, 0); + } + if (y >= polygon->Vertices[nextVR]->FinalPosition[1] && curVR != polygon->VBottom) + { + while (y >= polygon->Vertices[nextVR]->FinalPosition[1] && curVR != polygon->VBottom) + { + curVR = nextVR; + if (polygon->FacingView) + { + nextVR = curVR - 1; + if ((s32)nextVR < 0) + nextVR = nverts - 1; + } + else + { + nextVR = curVR + 1; + if (nextVR >= nverts) + nextVR = 0; + } + } + + if (polygon->Vertices[curVR]->FinalPosition[0] < minX) + { + minX = polygon->Vertices[curVR]->FinalPosition[0]; + minXY = polygon->Vertices[curVR]->FinalPosition[1]; + } + if (polygon->Vertices[curVR]->FinalPosition[0] > maxX) + { + maxX = polygon->Vertices[curVR]->FinalPosition[0]; + maxXY = polygon->Vertices[curVR]->FinalPosition[1]; + } + + assert(numYSpans < MaxYSpanSetups); + curSpanR = numYSpans; + SetupYSpan(i,&YSpanSetups[numYSpans++], polygon, curVR, nextVR, y, 1); + } + + assert(numSetupIndices < MaxYSpanIndices); + YSpanIndices[numSetupIndices].PolyIdx = i; + YSpanIndices[numSetupIndices].SpanIdxL = curSpanL; + YSpanIndices[numSetupIndices].SpanIdxR = curSpanR; + YSpanIndices[numSetupIndices].Y = y; + numSetupIndices++; + } + } + + if (polygon->Vertices[nextVL]->FinalPosition[0] < minX) + { + minX = polygon->Vertices[nextVL]->FinalPosition[0]; + minXY = polygon->Vertices[nextVL]->FinalPosition[1]; + } + if (polygon->Vertices[nextVL]->FinalPosition[0] > maxX) + { + maxX = polygon->Vertices[nextVL]->FinalPosition[0]; + maxXY = polygon->Vertices[nextVL]->FinalPosition[1]; + } + if (polygon->Vertices[nextVR]->FinalPosition[0] < minX) + { + minX = polygon->Vertices[nextVR]->FinalPosition[0]; + minXY = polygon->Vertices[nextVR]->FinalPosition[1]; + } + if (polygon->Vertices[nextVR]->FinalPosition[0] > maxX) + { + maxX = polygon->Vertices[nextVR]->FinalPosition[0]; + maxXY = polygon->Vertices[nextVR]->FinalPosition[1]; + } + + RenderPolygons[i].XMin = minX; + RenderPolygons[i].XMinY = minXY; + RenderPolygons[i].XMax = maxX; + RenderPolygons[i].XMaxY = maxXY; + + //printf("polygon min max %d %d | %d %d\n", RenderPolygons[i].XMin, RenderPolygons[i].XMinY, RenderPolygons[i].XMax, RenderPolygons[i].XMaxY); + } + + /*for (u32 i = 0; i < RenderNumPolygons; i++) + { + if (RenderPolygons[i].Variant >= numVariants) + { + printf("blarb2 %d %d %d\n", RenderPolygons[i].Variant, i, RenderNumPolygons); + } + //assert(RenderPolygons[i].Variant < numVariants); + }*/ + + if (numYSpans > 0) + { + glBindBuffer(GL_SHADER_STORAGE_BUFFER, YSpanSetupMemory); + glBufferSubData(GL_SHADER_STORAGE_BUFFER, 0, sizeof(SpanSetupY)*numYSpans, YSpanSetups); + + glBindBuffer(GL_TEXTURE_BUFFER, YSpanIndicesTextureMemory); + glBufferSubData(GL_TEXTURE_BUFFER, 0, numSetupIndices*4*2, YSpanIndices); + + glBindBuffer(GL_SHADER_STORAGE_BUFFER, RenderPolygonMemory); + glBufferSubData(GL_SHADER_STORAGE_BUFFER, 0, RenderNumPolygons*sizeof(RenderPolygon), RenderPolygons); + // we haven't accessed image data yet, so we don't need to invalidate anything + } + + //printf("found via %d %d %d of %d\n", foundviatexcache, foundviaprev, numslow, RenderNumPolygons); + + // bind everything + glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 0, YSpanSetupMemory); + glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, XSpanSetupMemory); + glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 2, RenderPolygonMemory); + glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 3, BinResultMemory); + glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 4, TileMemory); + glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 5, FinalTileMemory); + + MetaUniform meta; + meta.DispCnt = RenderDispCnt; + meta.NumPolygons = RenderNumPolygons; + meta.NumVariants = numVariants; + meta.AlphaRef = RenderAlphaRef; + { + u32 r = (RenderClearAttr1 << 1) & 0x3E; if (r) r++; + u32 g = (RenderClearAttr1 >> 4) & 0x3E; if (g) g++; + u32 b = (RenderClearAttr1 >> 9) & 0x3E; if (b) b++; + u32 a = (RenderClearAttr1 >> 16) & 0x1F; + meta.ClearColor = r | (g << 8) | (b << 16) | (a << 24); + meta.ClearDepth = ((RenderClearAttr2 & 0x7FFF) * 0x200) + 0x1FF; + meta.ClearAttr = RenderClearAttr1 & 0x3F008000; + } + for (u32 i = 0; i < 32; i++) + { + u32 color = RenderToonTable[i]; + u32 r = (color << 1) & 0x3E; + u32 g = (color >> 4) & 0x3E; + u32 b = (color >> 9) & 0x3E; + if (r) r++; + if (g) g++; + if (b) b++; + + meta.ToonTable[i*4+0] = r | (g << 8) | (b << 16); + } + for (u32 i = 0; i < 34; i++) + { + meta.ToonTable[i*4+1] = RenderFogDensityTable[i]; + } + for (u32 i = 0; i < 8; i++) + { + u32 color = RenderEdgeTable[i]; + u32 r = (color << 1) & 0x3E; + u32 g = (color >> 4) & 0x3E; + u32 b = (color >> 9) & 0x3E; + if (r) r++; + if (g) g++; + if (b) b++; + + meta.ToonTable[i*4+2] = r | (g << 8) | (b << 16); + } + meta.FogOffset = RenderFogOffset; + meta.FogShift = RenderFogShift; + { + u32 fogR = (RenderFogColor << 1) & 0x3E; if (fogR) fogR++; + u32 fogG = (RenderFogColor >> 4) & 0x3E; if (fogG) fogG++; + u32 fogB = (RenderFogColor >> 9) & 0x3E; if (fogB) fogB++; + u32 fogA = (RenderFogColor >> 16) & 0x1F; + meta.FogColor = fogR | (fogG << 8) | (fogB << 16) | (fogA << 24); + } + meta.XScroll = RenderXPos; + + glBindBuffer(GL_UNIFORM_BUFFER, MetaUniformMemory); + glBufferSubData(GL_UNIFORM_BUFFER, 0, sizeof(MetaUniform), &meta); + glBindBufferBase(GL_UNIFORM_BUFFER, 0, MetaUniformMemory); + + glUseProgram(ShaderClearCoarseBinMask); + glDispatchCompute(TilesPerLine*TileLines/32, 1, 1); + + bool wbuffer = false; + if (numYSpans > 0) + { + wbuffer = RenderPolygonRAM[0]->WBuffer; + + glUseProgram(ShaderClearIndirectWorkCount); + glDispatchCompute((numVariants+31)/32, 1, 1); + + // calculate x-spans + glBindImageTexture(0, YSpanIndicesTexture, 0, GL_FALSE, 0, GL_READ_ONLY, GL_RGBA16UI); + glUseProgram(ShaderInterpXSpans[wbuffer]); + glDispatchCompute((numSetupIndices + 31) / 32, 1, 1); + glMemoryBarrier(GL_SHADER_STORAGE_BUFFER); + + // bin polygons + glUseProgram(ShaderBinCombined); + glDispatchCompute(((RenderNumPolygons + 31) / 32), 256/CoarseTileW, 192/CoarseTileH); + glMemoryBarrier(GL_SHADER_STORAGE_BUFFER); + + // calculate list offsets + glUseProgram(ShaderCalculateWorkListOffset); + glDispatchCompute((numVariants + 31) / 32, 1, 1); + glMemoryBarrier(GL_SHADER_STORAGE_BUFFER); + + + // sort shader work + glUseProgram(ShaderSortWork); + glBindBuffer(GL_DISPATCH_INDIRECT_BUFFER, BinResultMemory); + glDispatchComputeIndirect(offsetof(BinResult, SortWorkWorkCount)); + glMemoryBarrier(GL_SHADER_STORAGE_BUFFER); + + // rasterise + { + bool highLightMode = RenderDispCnt & (1<<1); + + GLuint shadersNoTexture[] = + { + ShaderRasteriseNoTexture[wbuffer], + ShaderRasteriseNoTexture[wbuffer], + highLightMode + ? ShaderRasteriseNoTextureHighlight[wbuffer] + : ShaderRasteriseNoTextureToon[wbuffer], + ShaderRasteriseNoTexture[wbuffer], + ShaderRasteriseShadowMask[wbuffer] + }; + GLuint shadersUseTexture[] = + { + ShaderRasteriseUseTextureModulate[wbuffer], + ShaderRasteriseUseTextureDecal[wbuffer], + highLightMode + ? ShaderRasteriseUseTextureHighlight[wbuffer] + : ShaderRasteriseUseTextureToon[wbuffer], + ShaderRasteriseUseTextureDecal[wbuffer], + ShaderRasteriseShadowMask[wbuffer] + }; + + GLuint prevShader = 0; + s32 prevTexture = -1, prevSampler = -1; + for (int i = 0; i < numVariants; i++) + { + GLuint shader = 0; + if (variants[i].Texture == -1) + { + shader = shadersNoTexture[variants[i].BlendMode]; + } + else + { + shader = shadersUseTexture[variants[i].BlendMode]; + if (variants[i].Texture != prevTexture || variants[i].Sampler != prevSampler) + { + assert(variants[i].Sampler < 9); + glBindTexture(GL_TEXTURE_2D, variants[i].Texture); + prevTexture = variants[i].Texture; + prevSampler = variants[i].Sampler; + } + } + assert(shader != 0); + if (shader != prevShader) + { + glUseProgram(shader); + prevShader = shader; + } + + glUniform1i(UniformIdxCurVariant, i); + glUniform2f(UniformIdxTextureSize, 1.f / variants[i].Width, 1.f / variants[i].Height); + glBindBuffer(GL_DISPATCH_INDIRECT_BUFFER, BinResultMemory); + glDispatchComputeIndirect(offsetof(BinResult, VariantWorkCount) + i*4*4); + } + } + } + glMemoryBarrier(GL_SHADER_STORAGE_BARRIER_BIT); + + // compose final image + glUseProgram(ShaderDepthBlend[wbuffer]); + glDispatchCompute(256/8, 192/8, 1); + glMemoryBarrier(GL_SHADER_STORAGE_BARRIER_BIT); + + //glBindTexture(GL_TEXTURE_2D, ) + u32 finalPassShader = 0; + if (RenderDispCnt & (1<<4)) + finalPassShader |= 0x4; + if (RenderDispCnt & (1<<7)) + finalPassShader |= 0x2; + if (RenderDispCnt & (1<<5)) + finalPassShader |= 0x1; + + glUseProgram(ShaderFinalPass[finalPassShader]); + glDispatchCompute(256/32, 192, 1); + glMemoryBarrier(GL_SHADER_IMAGE_ACCESS_BARRIER_BIT); + + /*u64 starttime = armGetSystemTick(); + EmuQueue.waitIdle(); + printf("total time %f\n", armTicksToNs(armGetSystemTick()-starttime)*0.000001f);*/ + + /*for (u32 i = 0; i < RenderNumPolygons; i++) + { + if (RenderPolygons[i].Variant >= numVariants) + { + printf("blarb %d %d %d\n", RenderPolygons[i].Variant, i, RenderNumPolygons); + } + //assert(RenderPolygons[i].Variant < numVariants); + }*/ + + /*for (int i = 0; i < binresult->SortWorkWorkCount[0]*32; i++) + { + printf("sorted %x %x\n", binresult->SortedWork[i*2+0], binresult->SortedWork[i*2+1]); + }*/ +/* if (polygonvisible != -1) + { + SpanSetupX* xspans = Gfx::DataHeap->CpuAddr(XSpanSetupMemory); + printf("span result\n"); + Polygon* poly = RenderPolygonRAM[polygonvisible]; + u32 xspanoffset = RenderPolygons[polygonvisible].FirstXSpan; + for (u32 i = 0; i < (poly->YBottom - poly->YTop); i++) + { + printf("%d: %d - %d | %d %d | %d %d\n", i + poly->YTop, xspans[xspanoffset + i].X0, xspans[xspanoffset + i].X1, xspans[xspanoffset + i].__pad0, xspans[xspanoffset + i].__pad1, RenderPolygons[polygonvisible].YTop, RenderPolygons[polygonvisible].YBot); + } + }*/ +/* + printf("xspans: %d\n", numSetupIndices); + SpanSetupX* xspans = Gfx::DataHeap->CpuAddr(XSpanSetupMemory[curSlice]); + for (int i = 0; i < numSetupIndices; i++) + { + printf("poly %d %d %d | line %d | %d to %d\n", YSpanIndices[i].PolyIdx, YSpanIndices[i].SpanIdxL, YSpanIndices[i].SpanIdxR, YSpanIndices[i].Y, xspans[i].X0, xspans[i].X1); + } + printf("bin result\n"); + BinResult* binresult = Gfx::DataHeap->CpuAddr(BinResultMemory); + for (u32 y = 0; y < 192/8; y++) + { + for (u32 x = 0; x < 256/8; x++) + { + printf("%08x ", binresult->BinnedMaskCoarse[(x + y * (256/8)) * 2]); + } + printf("\n"); + }*/ +} + +void ComputeRenderer::RestartFrame() +{ + +} + +u32* ComputeRenderer::GetLine(int line) +{ + return DummyLine; +} + +} \ No newline at end of file diff --git a/src/GPU3D_Compute.h b/src/GPU3D_Compute.h new file mode 100644 index 00000000..5278c76c --- /dev/null +++ b/src/GPU3D_Compute.h @@ -0,0 +1,263 @@ +/* + Copyright 2016-2022 melonDS team + + This file is part of melonDS. + + melonDS is free software: you can redistribute it and/or modify it under + the terms of the GNU General Public License as published by the Free + Software Foundation, either version 3 of the License, or (at your option) + any later version. + + melonDS is distributed in the hope that it will be useful, but WITHOUT ANY + WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS + FOR A PARTICULAR PURPOSE. See the GNU General Public License for more details. + + You should have received a copy of the GNU General Public License along + with melonDS. If not, see http://www.gnu.org/licenses/. +*/ + +#ifndef GPU3D_COMPUTE +#define GPU3D_COMPUTE + +#include "GPU3D.h" + +#include "OpenGLSupport.h" + +#include "NonStupidBitfield.h" + +#include + +namespace GPU3D +{ + +class ComputeRenderer : public Renderer3D +{ +public: + ComputeRenderer(); + ~ComputeRenderer() override; + + bool Init() override; + void DeInit() override; + void Reset() override; + + void SetRenderSettings(GPU::RenderSettings& settings) override; + + void VCount144() override; + + void RenderFrame() override; + void RestartFrame() override; + u32* GetLine(int line) override; + + //dk::Fence FrameReady = {}; + //dk::Fence FrameReserveFence = {}; +private: + GLuint ShaderInterpXSpans[2]; + GLuint ShaderBinCombined; + GLuint ShaderDepthBlend[2]; + GLuint ShaderRasteriseNoTexture[2]; + GLuint ShaderRasteriseNoTextureToon[2]; + GLuint ShaderRasteriseNoTextureHighlight[2]; + GLuint ShaderRasteriseUseTextureDecal[2]; + GLuint ShaderRasteriseUseTextureModulate[2]; + GLuint ShaderRasteriseUseTextureToon[2]; + GLuint ShaderRasteriseUseTextureHighlight[2]; + GLuint ShaderRasteriseShadowMask[2]; + GLuint ShaderClearCoarseBinMask; + GLuint ShaderClearIndirectWorkCount; + GLuint ShaderCalculateWorkListOffset; + GLuint ShaderSortWork; + GLuint ShaderFinalPass[8]; + + GLuint YSpanIndicesTextureMemory; + GLuint YSpanIndicesTexture; + GLuint YSpanSetupMemory; + GLuint XSpanSetupMemory; + GLuint BinResultMemory; + GLuint RenderPolygonMemory; + GLuint TileMemory; + GLuint FinalTileMemory; + + u32 DummyLine[256] = {}; + + struct SpanSetupY + { + // Attributes + s32 Z0, Z1, W0, W1; + s32 ColorR0, ColorG0, ColorB0; + s32 ColorR1, ColorG1, ColorB1; + s32 TexcoordU0, TexcoordV0; + s32 TexcoordU1, TexcoordV1; + + // Interpolator + s32 I0, I1; + s32 Linear; + s32 IRecip; + s32 W0n, W0d, W1d; + + // Slope + s32 Increment; + + s32 X0, X1, Y0, Y1; + s32 XMin, XMax; + s32 DxInitial; + + s32 XCovIncr; + u32 IsDummy, __pad1; + }; + struct SpanSetupX + { + s32 X0, X1; + + s32 EdgeLenL, EdgeLenR, EdgeCovL, EdgeCovR; + + s32 XRecip; + + u32 Flags; + + s32 Z0, Z1, W0, W1; + s32 ColorR0, ColorG0, ColorB0; + s32 ColorR1, ColorG1, ColorB1; + s32 TexcoordU0, TexcoordV0; + s32 TexcoordU1, TexcoordV1; + + s32 CovLInitial, CovRInitial; + }; + struct SetupIndices + { + u16 PolyIdx, SpanIdxL, SpanIdxR, Y; + }; + struct RenderPolygon + { + u32 FirstXSpan; + s32 YTop, YBot; + + s32 XMin, XMax; + s32 XMinY, XMaxY; + + u32 Variant; + u32 Attr; + + float TextureLayer; + u32 __pad0, __pad1; + }; + + static const int TileSize = 8; + static const int CoarseTileCountX = 8; + static const int CoarseTileCountY = 4; + static const int CoarseTileW = CoarseTileCountX * TileSize; + static const int CoarseTileH = CoarseTileCountY * TileSize; + + static const int TilesPerLine = 256/TileSize; + static const int TileLines = 192/TileSize; + + static const int BinStride = 2048/32; + static const int CoarseBinStride = BinStride/32; + + static const int MaxWorkTiles = TilesPerLine*TileLines*48; + static const int MaxVariants = 256; + + static const int UniformIdxCurVariant = 0; + static const int UniformIdxTextureSize = 1; + + struct BinResult + { + u32 VariantWorkCount[MaxVariants*4]; + u32 SortedWorkOffset[MaxVariants]; + + u32 SortWorkWorkCount[4]; + u32 UnsortedWorkDescs[MaxWorkTiles*2]; + u32 SortedWork[MaxWorkTiles*2]; + + u32 BinnedMaskCoarse[TilesPerLine*TileLines*CoarseBinStride]; + u32 BinnedMask[TilesPerLine*TileLines*BinStride]; + u32 WorkOffsets[TilesPerLine*TileLines*BinStride]; + }; + + struct Tiles + { + u32 ColorTiles[MaxWorkTiles*TileSize*TileSize]; + u32 DepthTiles[MaxWorkTiles*TileSize*TileSize]; + u32 AttrStencilTiles[MaxWorkTiles*TileSize*TileSize]; + }; + + struct FinalTiles + { + u32 ColorResult[256*192*2]; + u32 DepthResult[256*192*2]; + u32 AttrResult[256*192*2]; + }; + + // eh those are pretty bad guesses + // though real hw shouldn't be eable to render all 2048 polygons on every line either + static const int MaxYSpanIndices = 64*2048; + static const int MaxYSpanSetups = 6144*2; + SetupIndices YSpanIndices[MaxYSpanIndices]; + SpanSetupY YSpanSetups[MaxYSpanSetups]; + RenderPolygon RenderPolygons[2048]; + + struct TexArrayEntry + { + u16 TexArrayIdx; + u16 LayerIdx; + }; + struct TexArray + { + GLuint Image; + u32 ImageDescriptor; + }; + + struct TexCacheEntry + { + u32 DescriptorIdx; + u32 LastVariant; // very cheap way to make variant lookup faster + + u32 TextureRAMStart[2], TextureRAMSize[2]; + u32 TexPalStart, TexPalSize; + u8 WidthLog2, HeightLog2; + TexArrayEntry Texture; + + u64 TextureHash[2]; + u64 TexPalHash; + }; + std::unordered_map TexCache; + + struct MetaUniform + { + u32 NumPolygons; + u32 NumVariants; + + u32 AlphaRef; + u32 DispCnt; + + u32 ToonTable[4*34]; + + u32 ClearColor, ClearDepth, ClearAttr; + + u32 FogOffset, FogShift, FogColor; + + u32 XScroll; + }; + GLuint MetaUniformMemory; + + static const u32 TexCacheMaxImages = 4096; + + u32 FreeImageDescriptorsCount = 0; + u32 FreeImageDescriptors[TexCacheMaxImages]; + + std::vector FreeTextures[8][8]; + std::vector TexArrays[8][8]; + + u32 TextureDecodingBuffer[1024*1024]; + + TexCacheEntry& GetTexture(u32 textureParam, u32 paletteParam); + + void SetupAttrs(SpanSetupY* span, Polygon* poly, int from, int to); + void SetupYSpan(int polynum, SpanSetupY* span, Polygon* poly, int from, int to, u32 y, int side); + void SetupYSpanDummy(SpanSetupY* span, Polygon* poly, int vertex, int side); + + bool CompileShader(GLuint& shader, const char* source, const std::initializer_list& defines); +}; + +} + +#endif \ No newline at end of file diff --git a/src/GPU3D_Compute_shaders.h b/src/GPU3D_Compute_shaders.h new file mode 100644 index 00000000..d770658c --- /dev/null +++ b/src/GPU3D_Compute_shaders.h @@ -0,0 +1,1502 @@ +/* + Copyright 2016-2022 melonDS team + + This file is part of melonDS. + + melonDS is free software: you can redistribute it and/or modify it under + the terms of the GNU General Public License as published by the Free + Software Foundation, either version 3 of the License, or (at your option) + any later version. + + melonDS is distributed in the hope that it will be useful, but WITHOUT ANY + WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS + FOR A PARTICULAR PURPOSE. See the GNU General Public License for more details. + + You should have received a copy of the GNU General Public License along + with melonDS. If not, see http://www.gnu.org/licenses/. +*/ + +#ifndef GPU3D_COMPUTE_SHADERS +#define GPU3D_COMPUTE_SHADERS + +namespace GPU3D +{ + +namespace ComputeRendererShaders +{ + +// defines: +// InterpSpans +// BinCombined +// Rasterise +// DepthBlend +// ClearCoarseBinMask +// ClearIndirectWorkCount +// CalculateWorkOffsets +// SortWork +// FinalPass + +// AntiAliasing +// EdgeMarking +// Fog + +// ZBuffer +// WBuffer + +// for Rasterise +// NoTexture +// UseTexture +// Decal +// Modulate +// Toon +// Highlight +// ShadowMask + +const char* Common = R"( +struct Polygon +{ + int FirstXSpan; + int YTop, YBot; + + int XMin, XMax; + int XMinY, XMaxY; + + int Variant; + + uint Attr; + + float TextureLayer; +}; + +struct YSpanSetup +{ + // Attributes + int Z0, Z1, W0, W1; + int ColorR0, ColorG0, ColorB0; + int ColorR1, ColorG1, ColorB1; + int TexcoordU0, TexcoordV0; + int TexcoordU1, TexcoordV1; + + // Interpolator + int I0, I1; + bool Linear; + int IRecip; + int W0n, W0d, W1d; + + // Slope + int Increment; + + int X0, X1, Y0, Y1; + int XMin, XMax; + int DxInitial; + + int XCovIncr; + + bool IsDummy; +}; + +const uint XSpanSetup_Linear = 1U << 0; +const uint XSpanSetup_FillInside = 1U << 1; +const uint XSpanSetup_FillLeft = 1U << 2; +const uint XSpanSetup_FillRight = 1U << 3; + +struct XSpanSetup +{ + int X0, X1; + + int InsideStart, InsideEnd, EdgeCovL, EdgeCovR; + + int XRecip; + + uint Flags; + + int Z0, Z1, W0, W1; + int ColorR0, ColorG0, ColorB0; + int ColorR1, ColorG1, ColorB1; + int TexcoordU0, TexcoordV0; + int TexcoordU1, TexcoordV1; + + int CovLInitial, CovRInitial; +}; + +layout (std140, binding = 0) readonly buffer YSpanSetupsBuffer +{ + YSpanSetup YSpanSetups[]; +}; + +#if defined(InterpSpans) || defined(BinCombined) || defined(Rasterise) +layout (std140, binding = 1) +#ifdef InterpSpans +writeonly +#endif +#if defined(BinCombined) || defined(Rasterise) +readonly +#endif +buffer XSpanSetupsBuffer +{ + XSpanSetup XSpanSetups[]; +}; +#endif + +layout (std140, binding = 2) readonly buffer PolygonBuffer +{ + Polygon Polygons[]; +}; + +#define TileSize 8 +const int CoarseTileCountX = 8; +const int CoarseTileCountY = 4; +const int CoarseTileW = (CoarseTileCountX * TileSize); +const int CoarseTileH = (CoarseTileCountY * TileSize); + +const int FramebufferStride = 256*192; +const int TilesPerLine = 256/TileSize; +const int TileLines = 192/TileSize; + +const int BinStride = 2048/32; +const int CoarseBinStride = BinStride/32; + +const int MaxWorkTiles = TilesPerLine*TileLines*48; +const int MaxVariants = 256; + +layout (std430, binding = 3) +buffer BinResultBuffer +{ + uvec4 VariantWorkCount[MaxVariants]; + uint SortedWorkOffset[MaxVariants]; + + uvec4 SortWorkWorkCount; + uvec2 UnsortedWorkDescs[MaxWorkTiles]; + uvec2 SortedWork[MaxWorkTiles]; + + uint BinnedMaskCoarse[TilesPerLine*TileLines*CoarseBinStride]; + uint BinnedMask[TilesPerLine*TileLines*BinStride]; + uint WorkOffsets[TilesPerLine*TileLines*BinStride]; +}; + +#if defined(Rasterise) || defined(DepthBlend) +layout (std430, binding = 4) +#ifdef Rasterise +writeonly +#endif +#ifdef DepthBlend +readonly +#endif +buffer TilesBuffer +{ + uint ColorTiles[MaxWorkTiles*TileSize*TileSize]; + uint DepthTiles[MaxWorkTiles*TileSize*TileSize]; + uint AttrTiles[MaxWorkTiles*TileSize*TileSize]; +}; +#endif + +layout (std430, binding = 5) +#ifdef DepthBlend +writeonly +#endif +#ifdef FinalPass +readonly +#endif +buffer RasterResult +{ + uint ColorResult[256*192*2]; + uint DepthResult[256*192*2]; + uint AttrResult[256*192*2]; +}; + +layout (std140, binding = 0) uniform MetaUniform +{ + uint NumPolygons; + uint NumVariants; + + int AlphaRef; + + uint DispCnt; + + // r = Toon + // g = Fog Density + // b = Edge Color + uvec4 ToonTable[34]; + + uint ClearColor, ClearDepth, ClearAttr; + + uint FogOffset, FogShift, FogColor; + + int XScroll; + + // only used/updated for rasteriation + uint CurVariant; + vec2 InvTextureSize; +}; + + +#if defined(InterpSpans) || defined(Rasterise) +uint Umulh(uint a, uint b) +{ + uint lo, hi; + umulExtended(a, b, hi, lo); + return hi; +} + +const uint startTable[256] = uint[256]( + 254, 252, 250, 248, 246, 244, 242, 240, 238, 236, 234, 233, 231, 229, 227, 225, 224, 222, 220, 218, 217, 215, 213, 212, 210, 208, 207, 205, 203, 202, 200, 199, 197, 195, 194, 192, 191, 189, 188, 186, 185, 183, 182, 180, 179, 178, 176, 175, 173, 172, 170, 169, 168, 166, 165, 164, 162, 161, 160, 158, +157, 156, 154, 153, 152, 151, 149, 148, 147, 146, 144, 143, 142, 141, 139, 138, 137, 136, 135, 134, 132, 131, 130, 129, 128, 127, 126, 125, 123, 122, 121, 120, 119, 118, 117, 116, 115, 114, 113, 112, 111, 110, 109, 108, 107, 106, 105, 104, 103, 102, 101, 100, 99, 98, 97, 96, 95, 94, 93, 92, 91, 90, 89, 88, 88, 87, 86, 85, 84, 83, 82, 81, 80, 80, 79, 78, 77, 76, 75, 74, 74, 73, 72, 71, 70, 70, 69, 68, 67, 66, 66, 65, 64, 63, 62, 62, 61, 60, 59, 59, 58, 57, 56, 56, 55, 54, 53, 53, 52, 51, 50, 50, 49, 48, 48, 47, 46, 46, 45, 44, 43, 43, 42, 41, 41, 40, 39, 39, 38, 37, 37, 36, 35, 35, 34, 33, 33, 32, 32, 31, 30, 30, 29, 28, 28, 27, 27, 26, 25, 25, 24, 24, 23, 22, 22, 21, 21, 20, 19, 19, 18, 18, 17, 17, 16, 15, 15, 14, 14, 13, 13, 12, 12, 11, 10, 10, 9, 9, 8, 8, 7, 7, 6, 6, 5, 5, 4, 4, 3, 3, 2, 2, 1, 1, 0, 0 +); + +uint Div(uint x, uint y) +{ + // https://www.microsoft.com/en-us/research/publication/software-integer-division/ + uint k = 31 - findMSB(y); + uint ty = (y << k) >> (32 - 9); + uint t = startTable[ty - 256] + 256; + uint z = (t << (32 - 9)) >> (32 - k - 1); + uint my = 0 - y; + + z += Umulh(z, my * z); + z += Umulh(z, my * z); + + uint q = Umulh(x, z); + uint r = x - y * q; + if(r >= y) + { + r = r - y; + q = q + 1; + if(r >= y) + { + r = r - y; + q = q + 1; + } + } + + return q; +} + +#ifdef InterpSpans +const int Shift = 9; +#else +const int Shift = 8; +#endif + +int CalcYFactorY(YSpanSetup span, int i) +{ + int num = abs((i) * span.W0n) << Shift; + int den = abs(((i) * span.W0d) + (((span.I1 - span.I0 - i) * span.W1d))); + + if (den == 0) + { + return 0; + } + else + { + int q = int(Div(num, den)); + //if ((num < 0) != (den < 0)) + // return -q; + return q; + } +} + +int CalcYFactorX(XSpanSetup span, int x) +{ + x -= span.X0; + + if (span.X0 != span.X1) + { + uint num = (uint(x) * span.W0) << Shift; + uint den = (uint(x) * span.W0) + (uint(span.X1 - span.X0 - x) * span.W1); + + if (den == 0) + return 0; + else + return int(Div(num, den)); + } + else + { + return 0; + } +} + +int InterpolateAttrPersp(int y0, int y1, int ifactor) +{ + if (y0 == y1) + return y0; + + if (y0 < y1) + return y0 + (((y1-y0) * ifactor) >> Shift); + else + return y1 + (((y0-y1) * ((1<> Shift); +} + +int InterpolateAttrLinear(int y0, int y1, int i, int irecip, int idiff) +{ + if (y0 == y1) + return y0; + +#ifndef Rasterise + irecip = abs(irecip); +#endif + + uint mulLo, mulHi, carry; + if (y0 < y1) + { +#ifndef Rasterise + uint offset = uint(abs(i)); +#else + uint offset = uint(i); +#endif + umulExtended(uint(y1-y0)*offset, uint(irecip), mulHi, mulLo); + mulLo = uaddCarry(mulLo, 3U<<24, carry); + mulHi += carry; + return y0 + int((mulLo >> 30) | (mulHi << (32 - 30))); + //return y0 + int(((int64_t(y1-y0) * int64_t(offset) * int64_t(irecip)) + int64_t(3<<24)) >> 30); + } + else + { +#ifndef Rasterise + uint offset = uint(abs(idiff-i)); +#else + uint offset = uint(idiff-i); +#endif + umulExtended(uint(y0-y1)*offset, uint(irecip), mulHi, mulLo); + mulLo = uaddCarry(mulLo, 3<<24, carry); + mulHi += carry; + return y1 + int((mulLo >> 30) | (mulHi << (32 - 30))); + //return y1 + int(((int64_t(y0-y1) * int64_t(offset) * int64_t(irecip)) + int64_t(3<<24)) >> 30); + } +} + +uint InterpolateZZBuffer(int z0, int z1, int i, int irecip, int idiff) +{ + if (z0 == z1) + return z0; + + uint base, disp, factor; + if (z0 < z1) + { + base = uint(z0); + disp = uint(z1 - z0); + factor = uint(abs(i)); + } + else + { + base = uint(z1); + disp = uint(z0 - z1), + factor = uint(abs(idiff - i)); + } + +#ifdef InterpSpans + int shiftl = 0; + const int shiftr = 22; + if (disp > 0x3FF) + { + shiftl = findMSB(disp) - 9; + disp >>= shiftl; + } +#else + disp >>= 9; + const int shiftl = 0; + const int shiftr = 13; +#endif + uint mulLo, mulHi; + + umulExtended(disp * factor, abs(irecip) >> 8, mulHi, mulLo); + + return base + (((mulLo >> shiftr) | (mulHi << (32 - shiftr))) << shiftl); +/* + int base, disp, factor; + if (z0 < z1) + { + base = z0; + disp = z1 - z0; + factor = i; + } + else + { + base = z1; + disp = z0 - z1, + factor = idiff - i; + } + +#ifdef InterpSpans + { + int shift = 0; + while (disp > 0x3FF) + { + disp >>= 1; + shift++; + } + + return base + int(((int64_t(disp) * int64_t(factor) * (int64_t(irecip) >> 8)) >> 22) << shift); + } +#else + { + disp >>= 9; + return base + int((int64_t(disp) * int64_t(factor) * (int64_t(irecip) >> 8)) >> 13); + } +#endif*/ +} + +uint InterpolateZWBuffer(int z0, int z1, int ifactor) +{ + if (z0 == z1) + return z0; + +#ifdef Rasterise + // since the precision along x spans is only 8 bit the result will always fit in 32-bit + if (z0 < z1) + { + return uint(z0) + (((z1-z0) * ifactor) >> Shift); + } + else + { + return uint(z1) + (((z0-z1) * ((1<> Shift); + } +#else + uint mulLo, mulHi; + if (z0 < z1) + { + umulExtended(z1-z0, ifactor, mulHi, mulLo); + // 64-bit shift + return uint(z0) + ((mulLo >> Shift) | (mulHi << (32-Shift))); + } + else + { + umulExtended(z0-z1, (1<> Shift) | (mulHi << (32-Shift))); + } +#endif + /*if (z0 < z1) + { + return uint(z0) + uint((int64_t(z1-z0) * int64_t(ifactor)) >> Shift); + } + else + { + return uint(z1) + uint((int64_t(z0-z1) * int64_t((1<> Shift); + }*/ +} + +int CalculateDx(int y, YSpanSetup span) +{ + return span.DxInitial + (y - span.Y0) * span.Increment; +} + +int CalculateX(int dx, YSpanSetup span) +{ + int x = span.X0; + if (span.X1 < span.X0) + x -= dx >> 18; + else + x += dx >> 18; + return clamp(x, span.XMin, span.XMax); +} + +void EdgeParams_XMajor(bool side, int dx, YSpanSetup span, out int edgelen, out int edgecov) +{ + bool negative = span.X1 < span.X0; + int len; + if (side != negative) + len = (dx >> 18) - ((dx-span.Increment) >> 18); + else + len = ((dx+span.Increment) >> 18) - (dx >> 18); + edgelen = len; + + int xlen = span.XMax + 1 - span.XMin; + int startx = dx >> 18; + if (negative) startx = xlen - startx; + if (side) startx = startx - len + 1; + + int startcov = int(Div(uint(((startx << 10) + 0x1FF) * (span.Y1 - span.Y0)), uint(xlen))); + edgecov = (1<<31) | ((startcov & 0x3FF) << 12) | (span.XCovIncr & 0x3FF); +} + +void EdgeParams_YMajor(bool side, int dx, YSpanSetup span, out int edgelen, out int edgecov) +{ + bool negative = span.X1 < span.X0; + edgelen = 1; + + if (span.Increment == 0) + { + edgecov = 31; + } + else + { + int cov = ((dx >> 9) + (span.Increment >> 10)) >> 4; + if ((cov >> 5) != (dx >> 18)) cov = 31; + cov &= 0x1F; + if (side == negative) cov = 0x1F - cov; + + edgecov = cov; + } +} +#endif + +)"; + +const char* InterpSpans = R"( + +layout (local_size_x = 32) in; + +layout (binding = 0, rgba16ui) uniform readonly uimageBuffer SetupIndices; + +void main() +{ + uvec4 setup = imageLoad(SetupIndices, int(gl_GlobalInvocationID.x)); + + YSpanSetup spanL = YSpanSetups[setup.y]; + YSpanSetup spanR = YSpanSetups[setup.z]; + + XSpanSetup xspan; + xspan.Flags = 0U; + + int y = int(setup.w); + + int dxl = CalculateDx(y, spanL); + int dxr = CalculateDx(y, spanR); + + int xl = CalculateX(dxl, spanL); + int xr = CalculateX(dxr, spanR); + + Polygon polygon = Polygons[setup.x]; + + int edgeLenL, edgeLenR; + + if (xl > xr) + { + YSpanSetup tmpSpan = spanL; + spanL = spanR; + spanR = tmpSpan; + + int tmp = xl; + xl = xr; + xr = tmp; + + EdgeParams_YMajor(false, dxr, spanL, edgeLenL, xspan.EdgeCovL); + EdgeParams_YMajor(true, dxl, spanR, edgeLenR, xspan.EdgeCovR); + } + else + { + // edges are the right way + if (spanL.Increment > 0x40000) + EdgeParams_XMajor(false, dxl, spanL, edgeLenL, xspan.EdgeCovL); + else + EdgeParams_YMajor(false, dxl, spanL, edgeLenL, xspan.EdgeCovL); + if (spanR.Increment > 0x40000) + EdgeParams_XMajor(true, dxr, spanR, edgeLenR, xspan.EdgeCovR); + else + EdgeParams_YMajor(true, dxr, spanR, edgeLenR, xspan.EdgeCovR); + } + + xspan.CovLInitial = (xspan.EdgeCovL >> 12) & 0x3FF; + if (xspan.CovLInitial == 0x3FF) + xspan.CovLInitial = 0; + xspan.CovRInitial = (xspan.EdgeCovR >> 12) & 0x3FF; + if (xspan.CovRInitial == 0x3FF) + xspan.CovRInitial = 0; + + xspan.X0 = xl; + xspan.X1 = xr + 1; + + uint polyalpha = ((polygon.Attr >> 16) & 0x1FU); + bool isWireframe = polyalpha == 0U; + + if (!isWireframe || (y == polygon.YTop || y == polygon.YBot - 1)) + xspan.Flags |= XSpanSetup_FillInside; + + xspan.InsideStart = xspan.X0 + edgeLenL; + if (xspan.InsideStart > xspan.X1) + xspan.InsideStart = xspan.X1; + xspan.InsideEnd = xspan.X1 - edgeLenR; + if (xspan.InsideEnd > xspan.X1) + xspan.InsideEnd = xspan.X1; + + bool isShadowMask = ((polygon.Attr & 0x3F000030U) == 0x00000030U); + bool fillAllEdges = /*polyalpha < 31*/true; + + if (fillAllEdges || spanL.X1 < spanL.X0 || spanL.Increment <= 0x40000) + xspan.Flags |= XSpanSetup_FillLeft; + if (fillAllEdges || (spanR.X1 >= spanR.X0 && spanR.Increment > 0x40000) || spanR.Increment == 0) + xspan.Flags |= XSpanSetup_FillRight; + + if (spanL.I0 == spanL.I1) + { + xspan.TexcoordU0 = spanL.TexcoordU0; + xspan.TexcoordV0 = spanL.TexcoordV0; + xspan.ColorR0 = spanL.ColorR0; + xspan.ColorG0 = spanL.ColorG0; + xspan.ColorB0 = spanL.ColorB0; + xspan.Z0 = spanL.Z0; + xspan.W0 = spanL.W0; + } + else + { + int i = (spanL.Increment > 0x40000 ? xl : y) - spanL.I0; + int ifactor = CalcYFactorY(spanL, i); + int idiff = spanL.I1 - spanL.I0; + +#ifdef ZBuffer + xspan.Z0 = int(InterpolateZZBuffer(spanL.Z0, spanL.Z1, i, spanL.IRecip, idiff)); +#endif +#ifdef WBuffer + xspan.Z0 = int(InterpolateZWBuffer(spanL.Z0, spanL.Z1, ifactor)); +#endif + + if (!spanL.Linear) + { + xspan.TexcoordU0 = InterpolateAttrPersp(spanL.TexcoordU0, spanL.TexcoordU1, ifactor); + xspan.TexcoordV0 = InterpolateAttrPersp(spanL.TexcoordV0, spanL.TexcoordV1, ifactor); + + xspan.ColorR0 = InterpolateAttrPersp(spanL.ColorR0, spanL.ColorR1, ifactor); + xspan.ColorG0 = InterpolateAttrPersp(spanL.ColorG0, spanL.ColorG1, ifactor); + xspan.ColorB0 = InterpolateAttrPersp(spanL.ColorB0, spanL.ColorB1, ifactor); + + xspan.W0 = InterpolateAttrPersp(spanL.W0, spanL.W1, ifactor); + } + else + { + xspan.TexcoordU0 = InterpolateAttrLinear(spanL.TexcoordU0, spanL.TexcoordU1, i, spanL.IRecip, idiff); + xspan.TexcoordV0 = InterpolateAttrLinear(spanL.TexcoordV0, spanL.TexcoordV1, i, spanL.IRecip, idiff); + + xspan.ColorR0 = InterpolateAttrLinear(spanL.ColorR0, spanL.ColorR1, i, spanL.IRecip, idiff); + xspan.ColorG0 = InterpolateAttrLinear(spanL.ColorG0, spanL.ColorG1, i, spanL.IRecip, idiff); + xspan.ColorB0 = InterpolateAttrLinear(spanL.ColorB0, spanL.ColorB1, i, spanL.IRecip, idiff); + + xspan.W0 = spanL.W0; // linear mode is only taken if W0 == W1 + } + } + + if (spanR.I0 == spanR.I1) + { + xspan.TexcoordU1 = spanR.TexcoordU0; + xspan.TexcoordV1 = spanR.TexcoordV0; + xspan.ColorR1 = spanR.ColorR0; + xspan.ColorG1 = spanR.ColorG0; + xspan.ColorB1 = spanR.ColorB0; + xspan.Z1 = spanR.Z0; + xspan.W1 = spanR.W0; + } + else + { + int i = (spanR.Increment > 0x40000 ? xr : y) - spanR.I0; + int ifactor = CalcYFactorY(spanR, i); + int idiff = spanR.I1 - spanR.I0; + + #ifdef ZBuffer + xspan.Z1 = int(InterpolateZZBuffer(spanR.Z0, spanR.Z1, i, spanR.IRecip, idiff)); + #endif + #ifdef WBuffer + xspan.Z1 = int(InterpolateZWBuffer(spanR.Z0, spanR.Z1, ifactor)); + #endif + + if (!spanR.Linear) + { + xspan.TexcoordU1 = InterpolateAttrPersp(spanR.TexcoordU0, spanR.TexcoordU1, ifactor); + xspan.TexcoordV1 = InterpolateAttrPersp(spanR.TexcoordV0, spanR.TexcoordV1, ifactor); + + xspan.ColorR1 = InterpolateAttrPersp(spanR.ColorR0, spanR.ColorR1, ifactor); + xspan.ColorG1 = InterpolateAttrPersp(spanR.ColorG0, spanR.ColorG1, ifactor); + xspan.ColorB1 = InterpolateAttrPersp(spanR.ColorB0, spanR.ColorB1, ifactor); + + xspan.W1 = int(InterpolateAttrPersp(spanR.W0, spanR.W1, ifactor)); + } + else + { + xspan.TexcoordU1 = InterpolateAttrLinear(spanR.TexcoordU0, spanR.TexcoordU1, i, spanR.IRecip, idiff); + xspan.TexcoordV1 = InterpolateAttrLinear(spanR.TexcoordV0, spanR.TexcoordV1, i, spanR.IRecip, idiff); + + xspan.ColorR1 = InterpolateAttrLinear(spanR.ColorR0, spanR.ColorR1, i, spanR.IRecip, idiff); + xspan.ColorG1 = InterpolateAttrLinear(spanR.ColorG0, spanR.ColorG1, i, spanR.IRecip, idiff); + xspan.ColorB1 = InterpolateAttrLinear(spanR.ColorB0, spanR.ColorB1, i, spanR.IRecip, idiff); + + xspan.W1 = spanR.W0; + } + } + + if (xspan.W0 == xspan.W1 && ((xspan.W0 | xspan.W1) & 0x7F) == 0) + { + xspan.Flags |= XSpanSetup_Linear; +// a bit hacky, but when wbuffering we only need to calculate xrecip for linear spans +#ifdef ZBuffer + } + { +#endif + xspan.XRecip = int(Div(1U<<30, uint(xspan.X1 - xspan.X0))); + } + + XSpanSetups[gl_GlobalInvocationID.x] = xspan; +} + +)"; + +const char* ClearIndirectWorkCount = R"( + +layout (local_size_x = 32) in; + +void main() +{ + VariantWorkCount[gl_GlobalInvocationID.x] = uvec4(1, 1, 0, 0); +} + +)"; + +const char* ClearCoarseBinMask = R"( + +layout (local_size_x = 32) in; + +void main() +{ + BinnedMaskCoarse[gl_GlobalInvocationID.x*CoarseBinStride+0] = 0; + BinnedMaskCoarse[gl_GlobalInvocationID.x*CoarseBinStride+1] = 0; +} + +)"; + +const char* BinCombined = R"( + +layout (local_size_x = 32) in; + +bool BinPolygon(Polygon polygon, ivec2 topLeft, ivec2 botRight) +{ + if (polygon.YTop > botRight.y || polygon.YBot <= topLeft.y) + return false; + + int polygonHeight = polygon.YBot - polygon.YTop; + + int polyInnerTopY = clamp(topLeft.y - polygon.YTop, 0, max(polygonHeight-1, 0)); + int polyInnerBotY = clamp(botRight.y - polygon.YTop, 0, max(polygonHeight-1, 0)); + + XSpanSetup xspanTop = XSpanSetups[polygon.FirstXSpan + polyInnerTopY]; + XSpanSetup xspanBot = XSpanSetups[polygon.FirstXSpan + polyInnerBotY]; + + int minXL; + if (polygon.XMinY >= topLeft.y && polygon.XMinY <= botRight.y) + minXL = polygon.XMin; + else + minXL = min(xspanTop.X0, xspanBot.X0); + + if (minXL > botRight.x) + return false; + + int maxXR; + if (polygon.XMaxY >= topLeft.y && polygon.XMaxY <= botRight.y) + maxXR = polygon.XMax; + else + maxXR = max(xspanTop.X1, xspanBot.X1) - 1; + + if (maxXR < topLeft.x) + return false; + + return true; +} + +shared uint mergedMaskShared; + +void main() +{ + int groupIdx = int(gl_WorkGroupID.x); + ivec2 coarseTile = ivec2(gl_WorkGroupID.yz); + +#if 0 + int localIdx = int(gl_SubGroupInvocationARB); +#else + int localIdx = int(gl_LocalInvocationIndex); + + if (localIdx == 0) + mergedMaskShared = 0U; + barrier(); +#endif + + int polygonIdx = groupIdx * 32 + localIdx; + + ivec2 coarseTopLeft = coarseTile * ivec2(CoarseTileW, CoarseTileH); + ivec2 coarseBotRight = coarseTopLeft + ivec2(CoarseTileW-1, CoarseTileH-1); + + bool binned = false; + if (polygonIdx < NumPolygons) + { + binned = BinPolygon(Polygons[polygonIdx], coarseTopLeft, coarseBotRight); + } + +#if 0 + uint mergedMask = unpackUint2x32(ballotARB(binned)).x; +#else + if (binned) + atomicOr(mergedMaskShared, 1U << localIdx); + barrier(); + uint mergedMask = mergedMaskShared; +#endif + + ivec2 fineTile = ivec2(localIdx & 0x7, localIdx >> 3); + + ivec2 fineTileTopLeft = coarseTopLeft + fineTile * ivec2(TileSize, TileSize); + ivec2 fineTileBotRight = fineTileTopLeft + ivec2(TileSize-1, TileSize-1); + + uint binnedMask = 0U; + while (mergedMask != 0U) + { + int bit = findLSB(mergedMask); + mergedMask &= ~(1U << bit); + + int polygonIdx = groupIdx * 32 + bit; + + if (BinPolygon(Polygons[polygonIdx], fineTileTopLeft, fineTileBotRight)) + binnedMask |= 1U << bit; + } + + int linearTile = fineTile.x + fineTile.y * TilesPerLine + coarseTile.x * CoarseTileCountX + coarseTile.y * TilesPerLine * CoarseTileCountY; + + BinnedMask[linearTile * BinStride + groupIdx] = binnedMask; + int coarseMaskIdx = linearTile * CoarseBinStride + (groupIdx >> 5); + if (binnedMask != 0U) + atomicOr(BinnedMaskCoarse[coarseMaskIdx], 1U << (groupIdx & 0x1F)); + + if (binnedMask != 0U) + { + uint workOffset = atomicAdd(VariantWorkCount[0].w, uint(bitCount(binnedMask))); + WorkOffsets[linearTile * BinStride + groupIdx] = workOffset; + + uint tilePositionCombined = bitfieldInsert(fineTileTopLeft.x, fineTileTopLeft.y, 16, 16); + + int idx = 0; + while (binnedMask != 0U) + { + int bit = findLSB(binnedMask); + binnedMask &= ~(1U << bit); + + int polygonIdx = groupIdx * 32 + bit; + int variantIdx = Polygons[polygonIdx].Variant; + + int inVariantOffset = int(atomicAdd(VariantWorkCount[variantIdx].z, 1)); + UnsortedWorkDescs[workOffset + idx] = uvec2(tilePositionCombined, bitfieldInsert(inVariantOffset, polygonIdx, 16, 16)); + + idx++; + } + } +} + +)"; + +const char* CalcOffsets = R"( + +layout (local_size_x = 32) in; + +void main() +{ + if (gl_GlobalInvocationID.x < NumVariants) + { + if (gl_GlobalInvocationID.x == 0) + { + // a bit of a cheat putting this here, but this shader won't run that often + SortWorkWorkCount = uvec4((VariantWorkCount[0].w + 31) / 32, 1, 1, 0); + } + SortedWorkOffset[gl_GlobalInvocationID.x] = atomicAdd(VariantWorkCount[1].w, VariantWorkCount[gl_GlobalInvocationID.x].z); + } +} + + +)"; + +const char* SortWork = R"( + +layout (local_size_x = 32) in; + +void main() +{ + if (gl_GlobalInvocationID.x < VariantWorkCount[0].w) + { + uvec2 workDesc = UnsortedWorkDescs[gl_GlobalInvocationID.x]; + int inVariantOffset = int(bitfieldExtract(workDesc.y, 0, 16)); + int polygonIdx = int(bitfieldExtract(workDesc.y, 16, 16)); + int variantIdx = Polygons[polygonIdx].Variant; + + int sortedIndex = int(SortedWorkOffset[variantIdx]) + inVariantOffset; + SortedWork[sortedIndex] = uvec2(workDesc.x, bitfieldInsert(workDesc.y, gl_GlobalInvocationID.x, 0, 16)); + } +} + +)"; + +const char* Rasterise = R"( + +layout (local_size_x = TileSize, local_size_y = TileSize) in; + +layout (binding = 0) uniform usampler2DArray CurrentTexture; + +void main() +{ + uvec2 workDesc = SortedWork[SortedWorkOffset[CurVariant] + gl_WorkGroupID.z]; + Polygon polygon = Polygons[bitfieldExtract(workDesc.y, 16, 16)]; + ivec2 position = ivec2(bitfieldExtract(workDesc.x, 0, 16), bitfieldExtract(workDesc.x, 16, 16)) + ivec2(gl_LocalInvocationID.xy); + int tileOffset = int(bitfieldExtract(workDesc.y, 0, 16)) * TileSize * TileSize + TileSize * int(gl_LocalInvocationID.y) + int(gl_LocalInvocationID.x); + + uint color = 0U; + if (position.y >= polygon.YTop && position.y < polygon.YBot) + { + XSpanSetup xspan = XSpanSetups[polygon.FirstXSpan + (position.y - polygon.YTop)]; + + bool insideLeftEdge = position.x < xspan.InsideStart; + bool insideRightEdge = position.x >= xspan.InsideEnd; + bool insidePolygonInside = !insideLeftEdge && !insideRightEdge; + + if (position.x >= xspan.X0 && position.x < xspan.X1 + && ((insideLeftEdge && (xspan.Flags & XSpanSetup_FillLeft) != 0U) + || (insideRightEdge && (xspan.Flags & XSpanSetup_FillRight) != 0U) + || (insidePolygonInside && (xspan.Flags & XSpanSetup_FillInside) != 0U))) + { + uint attr = 0; + if (position.y == polygon.YTop) + attr |= 0x4U; + else if (position.y == polygon.YBot - 1) + attr |= 0x8U; + + if (insideLeftEdge) + { + attr |= 0x1U; + + int cov = xspan.EdgeCovL; + if ((cov & (1U<<31)) != 0U) + { + int xcov = xspan.CovLInitial + (xspan.EdgeCovL & 0x3FF) * (position.x - xspan.X0); + cov = min(xcov >> 5, 31); + } + + attr |= uint(cov) << 8; + } + else if (insideRightEdge) + { + attr |= 0x2U; + + int cov = xspan.EdgeCovR; + if ((cov & (1U<<31)) != 0U) + { + int xcov = xspan.CovRInitial + (xspan.EdgeCovR & 0x3FF) * (position.x - xspan.InsideEnd); + cov = max(0x1F - (xcov >> 5), 0); + } + + attr |= uint(cov) << 8; + } + + uint z; + int u, v, vr, vg, vb; + + if (xspan.X0 == xspan.X1) + { + z = xspan.Z0; + u = xspan.TexcoordU0; + v = xspan.TexcoordV0; + vr = xspan.ColorR0; + vg = xspan.ColorG0; + vb = xspan.ColorB0; + } + else + { + int ifactor = CalcYFactorX(xspan, position.x); + int idiff = xspan.X1 - xspan.X0; + int i = position.x - xspan.X0; + +#ifdef ZBuffer + z = InterpolateZZBuffer(xspan.Z0, xspan.Z1, i, xspan.XRecip, idiff); +#endif +#ifdef WBuffer + z = InterpolateZWBuffer(xspan.Z0, xspan.Z1, ifactor); +#endif + if ((xspan.Flags & XSpanSetup_Linear) == 0U) + { + u = InterpolateAttrPersp(xspan.TexcoordU0, xspan.TexcoordU1, ifactor); + v = InterpolateAttrPersp(xspan.TexcoordV0, xspan.TexcoordV1, ifactor); + + vr = InterpolateAttrPersp(xspan.ColorR0, xspan.ColorR1, ifactor); + vg = InterpolateAttrPersp(xspan.ColorG0, xspan.ColorG1, ifactor); + vb = InterpolateAttrPersp(xspan.ColorB0, xspan.ColorB1, ifactor); + } + else + { + u = InterpolateAttrLinear(xspan.TexcoordU0, xspan.TexcoordU1, i, xspan.XRecip, idiff); + v = InterpolateAttrLinear(xspan.TexcoordV0, xspan.TexcoordV1, i, xspan.XRecip, idiff); + + vr = InterpolateAttrLinear(xspan.ColorR0, xspan.ColorR1, i, xspan.XRecip, idiff); + vg = InterpolateAttrLinear(xspan.ColorG0, xspan.ColorG1, i, xspan.XRecip, idiff); + vb = InterpolateAttrLinear(xspan.ColorB0, xspan.ColorB1, i, xspan.XRecip, idiff); + } + } + +#ifndef ShadowMask + vr >>= 3; + vg >>= 3; + vb >>= 3; + + uint r, g, b, a; + uint polyalpha = bitfieldExtract(polygon.Attr, 16, 5); + +#ifdef Toon + uint tooncolor = ToonTable[vr >> 1].r; + vr = int(bitfieldExtract(tooncolor, 0, 8)); + vg = int(bitfieldExtract(tooncolor, 8, 8)); + vb = int(bitfieldExtract(tooncolor, 16, 8)); +#endif +#ifdef Highlight + vg = vr; + vb = vr; +#endif + +#ifdef NoTexture + a = int(polyalpha); +#endif + r = vr; + g = vg; + b = vb; + +#ifdef UseTexture + vec2 uvf = vec2(ivec2(u, v)) * vec2(1.0 / 16.0) * InvTextureSize; + + uvec4 texcolor = texture(CurrentTexture, vec3(uvf, polygon.TextureLayer)); +#ifdef Decal + if (texcolor.a == 31) + { + r = int(texcolor.r); + g = int(texcolor.g); + b = int(texcolor.b); + } + else if (texcolor.a > 0) + { + r = int((texcolor.r * texcolor.a) + (vr * (31-texcolor.a))) >> 5; + g = int((texcolor.g * texcolor.a) + (vg * (31-texcolor.a))) >> 5; + b = int((texcolor.b * texcolor.a) + (vb * (31-texcolor.a))) >> 5; + } + a = int(polyalpha); +#endif +#if defined(Modulate) || defined(Toon) || defined(Highlight) + r = int((texcolor.r+1) * (vr+1) - 1) >> 6; + g = int((texcolor.g+1) * (vg+1) - 1) >> 6; + b = int((texcolor.b+1) * (vb+1) - 1) >> 6; + a = int((texcolor.a+1) * (polyalpha+1) - 1) >> 5; +#endif +#endif + +#ifdef Highlight + uint tooncolor = ToonTable[vr >> 1].r; + + r = min(r + int(bitfieldExtract(tooncolor, 0, 8)), 63); + g = min(g + int(bitfieldExtract(tooncolor, 8, 8)), 63); + b = min(b + int(bitfieldExtract(tooncolor, 16, 8)), 63); +#endif + + if (polyalpha == 0) + a = 31; + + if (a > AlphaRef) + { + color = r | (g << 8) | (b << 16) | (a << 24); + + DepthTiles[tileOffset] = z; + AttrTiles[tileOffset] = attr; + } +#else + color = 0xFFFFFFFF; // doesn't really matter as long as it's not 0 + DepthTiles[tileOffset] = z; +#endif + } + } + + ColorTiles[tileOffset] = color; +} + +)"; + +const char* DepthBlend = R"( + +layout (local_size_x = TileSize, local_size_y = TileSize) in; + +void PlotTranslucent(inout uint color, inout uint depth, inout uint attr, bool isShadow, uint tileColor, uint srcA, uint tileDepth, uint srcAttr, bool writeDepth) +{ + uint blendAttr = (srcAttr & 0xE0F0U) | ((srcAttr >> 8) & 0xFF0000U) | (1U<<22) | (attr & 0xFF001F0FU); + + if ((!isShadow || (attr & (1U<<22)) != 0U) + ? (attr & 0x007F0000U) != (blendAttr & 0x007F0000U) + : (attr & 0x3F000000U) != (srcAttr & 0x3F000000U)) + { + // le blend + if (writeDepth) + depth = tileDepth; + + if ((attr & (1U<<15)) == 0) + blendAttr &= ~(1U<<15); + attr = blendAttr; + + uint srcRB = tileColor & 0x3F003FU; + uint srcG = tileColor & 0x003F00U; + uint dstRB = color & 0x3F003FU; + uint dstG = color & 0x003F00U; + uint dstA = color & 0x1F000000U; + + uint alpha = (srcA >> 24) + 1; + if (dstA != 0) + { + srcRB = ((srcRB * alpha) + (dstRB * (32-alpha))) >> 5; + srcG = ((srcG * alpha) + (dstG * (32-alpha))) >> 5; + } + + color = (srcRB & 0x3F003FU) | (srcG & 0x003F00U) | max(dstA, srcA); + } +} + +void ProcessCoarseMask(int linearTile, uint coarseMask, uint coarseOffset, + inout uvec2 color, inout uvec2 depth, inout uvec2 attr, inout uint stencil, + inout bool prevIsShadowMask) +{ + int tileInnerOffset = int(gl_LocalInvocationID.x) + int(gl_LocalInvocationID.y) * TileSize; + + while (coarseMask != 0U) + { + uint coarseBit = findLSB(coarseMask); + coarseMask &= ~(1U << coarseBit); + + uint tileOffset = linearTile * BinStride + coarseBit + coarseOffset; + + uint fineMask = BinnedMask[tileOffset]; + uint workIdx = WorkOffsets[tileOffset]; + + while (fineMask != 0U) + { + uint fineIdx = findLSB(fineMask); + fineMask &= ~(1U << fineIdx); + + uint pixelindex = tileInnerOffset + workIdx * TileSize * TileSize; + uint tileColor = ColorTiles[pixelindex]; + workIdx++; + + uint polygonIdx = fineIdx + (coarseBit + coarseOffset) * 32; + + if (tileColor != 0U) + { + uint polygonAttr = Polygons[polygonIdx].Attr; + + bool isShadowMask = ((polygonAttr & 0x3F000030U) == 0x00000030U); + bool prevIsShadowMaskOld = prevIsShadowMask; + prevIsShadowMask = isShadowMask; + + bool equalDepthTest = (polygonAttr & (1U << 14)) != 0U; + + uint tileDepth = DepthTiles[pixelindex]; + uint tileAttr = AttrTiles[pixelindex]; + + uint dstattr = attr.x; + + if (!isShadowMask) + { + bool isShadow = (polygonAttr & 0x30U) == 0x30U; + + bool writeSecondLayer = false; + + if (isShadow) + { + if (stencil == 0U) + continue; + if ((stencil & 1U) == 0U) + writeSecondLayer = true; + if ((stencil & 2U) == 0U) + dstattr &= ~0x3U; + } + + uint dstDepth = writeSecondLayer ? depth.y : depth.x; + if (!(equalDepthTest +#ifdef WBuffer + ? dstDepth - tileDepth + 0xFFU <= 0x1FE +#endif +#ifdef ZBuffer + ? dstDepth - tileDepth + 0x200 <= 0x400 +#endif + : tileDepth < dstDepth)) + { + if ((dstattr & 0x3U) == 0U || writeSecondLayer) + continue; + + writeSecondLayer = true; + dstattr = attr.y; + if (!(equalDepthTest +#ifdef WBuffer + ? depth.y - tileDepth + 0xFFU <= 0x1FE +#endif +#ifdef ZBuffer + ? depth.y - tileDepth + 0x200 <= 0x400 +#endif + : tileDepth < depth.y)) + continue; + } + + uint srcAttr = (polygonAttr & 0x3F008000U); + + uint srcA = tileColor & 0x1F000000U; + if (srcA == 0x1F000000U) + { + srcAttr |= tileAttr; + + if (!writeSecondLayer) + { + if ((srcAttr & 0x3U) != 0U) + { + color.y = color.x; + depth.y = depth.x; + attr.y = attr.x; + } + + color.x = tileColor; + depth.x = tileDepth; + attr.x = srcAttr; + } + else + { + color.y = tileColor; + depth.y = tileDepth; + attr.y = srcAttr; + } + } + else + { + bool writeDepth = (polygonAttr & (1U<<11)) != 0; + + if (!writeSecondLayer) + { + // blend into both layers + PlotTranslucent(color.x, depth.x, attr.x, isShadow, tileColor, srcA, tileDepth, srcAttr, writeDepth); + } + if (writeSecondLayer || (dstattr & 0x3U) != 0U) + { + PlotTranslucent(color.y, depth.y, attr.y, isShadow, tileColor, srcA, tileDepth, srcAttr, writeDepth); + } + } + } + else + { + if (!prevIsShadowMaskOld) + stencil = 0; + + if (!(equalDepthTest +#ifdef WBuffer + ? depth.x - tileDepth + 0xFFU <= 0x1FE +#endif +#ifdef ZBuffer + ? depth.x - tileDepth + 0x200 <= 0x400 +#endif + : tileDepth < depth.x)) + stencil = 0x1U; + + if ((dstattr & 0x3U) != 0U) + { + if (!(equalDepthTest +#ifdef WBuffer + ? depth.y - tileDepth + 0xFFU <= 0x1FE +#endif +#ifdef ZBuffer + ? depth.y - tileDepth + 0x200 <= 0x400 +#endif + : tileDepth < depth.y)) + stencil |= 0x2U; + } + } + } + } + } +} + +void main() +{ + int linearTile = int(gl_WorkGroupID.x + (gl_WorkGroupID.y * TilesPerLine)); + + uint coarseMaskLo = BinnedMaskCoarse[linearTile*CoarseBinStride + 0]; + uint coarseMaskHi = BinnedMaskCoarse[linearTile*CoarseBinStride + 1]; + + uvec2 color = uvec2(ClearColor, 0U); + uvec2 depth = uvec2(ClearDepth, 0U); + uvec2 attr = uvec2(ClearAttr, 0U); + uint stencil = 0U; + bool prevIsShadowMask = false; + + ProcessCoarseMask(linearTile, coarseMaskLo, 0, color, depth, attr, stencil, prevIsShadowMask); + ProcessCoarseMask(linearTile, coarseMaskHi, BinStride/2, color, depth, attr, stencil, prevIsShadowMask); + + int resultOffset = int(gl_GlobalInvocationID.x) + int(gl_GlobalInvocationID.y) * 256; + ColorResult[resultOffset] = color.x; + ColorResult[resultOffset+FramebufferStride] = color.y; + DepthResult[resultOffset] = depth.x; + DepthResult[resultOffset+FramebufferStride] = depth.y; + AttrResult[resultOffset] = attr.x; + AttrResult[resultOffset+FramebufferStride] = attr.y; +} + +)"; + +const char* FinalPass = R"( + +layout (local_size_x = 32) in; + +layout (binding = 0, r32ui) writeonly uniform uimage2D FinalFB; + +uint BlendFog(uint color, uint depth) +{ + uint densityid = 0, densityfrac = 0; + + if (depth >= FogOffset) + { + depth -= FogOffset; + depth = (depth >> 2) << FogShift; + + densityid = depth >> 17; + if (densityid >= 32) + { + densityid = 32; + densityfrac = 0; + } + else + { + densityfrac = depth & 0x1FFFFU; + } + } + + uint density = + ((ToonTable[densityid].g * (0x20000U-densityfrac)) + + (ToonTable[densityid+1].g * densityfrac)) >> 17; + density = min(density, 128U); + + uint colorRB = color & 0x3F003FU; + uint colorGA = (color >> 8) & 0x3F003FU; + + uint fogRB = FogColor & 0x3F003FU; + uint fogGA = (FogColor >> 8) & 0x1F003FU; + + uint finalColorRB = ((fogRB * density) + (colorRB * (128-density))) >> 7; + uint finalColorGA = ((fogGA * density) + (colorGA * (128-density))) >> 7; + + finalColorRB &= 0x3F003FU; + finalColorGA &= 0x1F003FU; + + return (DispCnt & (1U<<6)) != 0 + ? (bitfieldInsert(color, finalColorGA >> 16, 24, 8)) + : (finalColorRB | (finalColorGA << 8)); +} + +void main() +{ + int srcX = (int(gl_GlobalInvocationID.x) + XScroll) & 0x1FF; + int resultOffset = int(srcX) + int(gl_GlobalInvocationID.y) * 256; + + uvec2 color = uvec2(0); + uvec2 depth = uvec2(0); + uvec2 attr = uvec2(0); + if (srcX < 256) + { + color = uvec2(ColorResult[resultOffset], ColorResult[resultOffset+FramebufferStride]); + depth = uvec2(DepthResult[resultOffset], DepthResult[resultOffset+FramebufferStride]); + attr = uvec2(AttrResult[resultOffset], AttrResult[resultOffset+FramebufferStride]); + } + +#ifdef EdgeMarking + if ((attr.x & 0xFU) != 0U) + { + uvec4 otherAttr = uvec4(ClearAttr); + uvec4 otherDepth = uvec4(ClearDepth); + + if (srcX > 0U) + { + otherAttr.x = AttrResult[resultOffset-1]; + otherDepth.x = DepthResult[resultOffset-1]; + } + if (srcX < 255U) + { + otherAttr.y = AttrResult[resultOffset+1]; + otherDepth.y = DepthResult[resultOffset+1]; + } + if (gl_GlobalInvocationID.y > 0U) + { + otherAttr.z = AttrResult[resultOffset-256]; + otherDepth.z = DepthResult[resultOffset-256]; + } + if (gl_GlobalInvocationID.y < 191U) + { + otherAttr.w = AttrResult[resultOffset+256]; + otherDepth.w = DepthResult[resultOffset+256]; + } + + uint polyId = bitfieldExtract(attr.x, 24, 5); + uvec4 otherPolyId = bitfieldExtract(otherAttr, 24, 5); + + bvec4 polyIdMatch = equal(uvec4(polyId), otherPolyId); + bvec4 nearer = lessThan(uvec4(depth.x), otherDepth); + + if ((!polyIdMatch.x && nearer.x) + || (!polyIdMatch.y && nearer.y) + || (!polyIdMatch.z && nearer.z) + || (!polyIdMatch.w && nearer.w)) + { + color.x = ToonTable[polyId >> 3].b | (color.x & 0xFF000000U); + attr.x = (attr.x & 0xFFFFE0FFU) | 0x00001000U; + } + } +#endif + +#ifdef Fog + if ((attr.x & (1U<<15)) != 0U) + { + color.x = BlendFog(color.x, depth.x); + } + + if ((attr.x & 0xFU) != 0 && (attr.y & (1U<<15)) != 0U) + { + color.y = BlendFog(color.y, depth.y); + } +#endif + +#ifdef AntiAliasing + // resolve anti-aliasing + if ((attr.x & 0x3U) != 0) + { + uint coverage = (attr.x >> 8) & 0x1FU; + + if (coverage != 0) + { + uint topRB = color.x & 0x3F003FU; + uint topG = color.x & 0x003F00U; + uint topA = bitfieldExtract(color.x, 24, 5); + + uint botRB = color.y & 0x3F003FU; + uint botG = color.y & 0x003F00U; + uint botA = bitfieldExtract(color.y, 24, 5); + + coverage++; + + if (botA > 0) + { + topRB = ((topRB * coverage) + (botRB * (32-coverage))) >> 5; + topG = ((topG * coverage) + (botG * (32-coverage))) >> 5; + + topRB &= 0x3F003FU; + topG &= 0x003F00U; + } + + topA = ((topA * coverage) + (botA * (32-coverage))) >> 5; + + color.x = topRB | topG | (topA << 24); + } + else + { + color.x = color.y; + } + } +#endif + + if (bitfieldExtract(color.x, 24, 8) != 0U) + color.x |= 0x40000000U; + else + color.x = 0U; + + //if (gl_LocalInvocationID.x == 7 || gl_LocalInvocationID.y == 7) + //color.x = 0x1F00001FU | 0x40000000U; + + imageStore(FinalFB, ivec2(gl_GlobalInvocationID.xy), uvec4(color.x, 0, 0, 0)); +} + +)"; + +} + +} + +#endif \ No newline at end of file diff --git a/src/GPU3D_OpenGL.cpp b/src/GPU3D_OpenGL.cpp index 9648be36..1e61dac2 100644 --- a/src/GPU3D_OpenGL.cpp +++ b/src/GPU3D_OpenGL.cpp @@ -49,25 +49,18 @@ bool GLRenderer::BuildRenderShader(u32 flags, const char* vs, const char* fs) strcpy(&fsbuf[headerlen], kRenderFSCommon); strcpy(&fsbuf[headerlen + fsclen], fs); - bool ret = OpenGL::BuildShaderProgram(vsbuf, fsbuf, RenderShader[flags], shadername); + GLuint prog; + bool ret = OpenGL::CompileVertexFragmentProgram(prog, + vsbuf, fsbuf, + shadername, + {{"vPosition", 0}, {"vColor", 1}, {"vTexcoord", 2}, {"vPolygonAttr", 3}}, + {{"oColor", 0}, {"oAttr", 1}}); delete[] vsbuf; delete[] fsbuf; if (!ret) return false; - GLuint prog = RenderShader[flags][2]; - - glBindAttribLocation(prog, 0, "vPosition"); - glBindAttribLocation(prog, 1, "vColor"); - glBindAttribLocation(prog, 2, "vTexcoord"); - glBindAttribLocation(prog, 3, "vPolygonAttr"); - glBindFragDataLocation(prog, 0, "oColor"); - glBindFragDataLocation(prog, 1, "oAttr"); - - if (!OpenGL::LinkShaderProgram(RenderShader[flags])) - return false; - GLint uni_id = glGetUniformBlockIndex(prog, "uConfig"); glUniformBlockBinding(prog, uni_id, 0); @@ -78,13 +71,15 @@ bool GLRenderer::BuildRenderShader(u32 flags, const char* vs, const char* fs) uni_id = glGetUniformLocation(prog, "TexPalMem"); glUniform1i(uni_id, 1); + RenderShader[flags] = prog; + return true; } void GLRenderer::UseRenderShader(u32 flags) { if (CurShaderID == flags) return; - glUseProgram(RenderShader[flags][2]); + glUseProgram(RenderShader[flags]); CurShaderID = flags; } @@ -118,21 +113,17 @@ std::unique_ptr GLRenderer::New() noexcept glDepthRange(0, 1); glClearDepth(1.0); - - if (!OpenGL::BuildShaderProgram(kClearVS, kClearFS, result->ClearShaderPlain, "ClearShader")) + if (!OpenGL::CompileVertexFragmentProgram(result->ClearShaderPlain, + kClearVS, kClearFS, + "ClearShader", + {{"vPosition", 0}}, + {{"oColor", 0}, {"oAttr", 1}})) return nullptr; - glBindAttribLocation(result->ClearShaderPlain[2], 0, "vPosition"); - glBindFragDataLocation(result->ClearShaderPlain[2], 0, "oColor"); - glBindFragDataLocation(result->ClearShaderPlain[2], 1, "oAttr"); - - if (!OpenGL::LinkShaderProgram(result->ClearShaderPlain)) - return nullptr; - - result->ClearUniformLoc[0] = glGetUniformLocation(result->ClearShaderPlain[2], "uColor"); - result->ClearUniformLoc[1] = glGetUniformLocation(result->ClearShaderPlain[2], "uDepth"); - result->ClearUniformLoc[2] = glGetUniformLocation(result->ClearShaderPlain[2], "uOpaquePolyID"); - result->ClearUniformLoc[3] = glGetUniformLocation(result->ClearShaderPlain[2], "uFogFlag"); + result->ClearUniformLoc[0] = glGetUniformLocation(result->ClearShaderPlain, "uColor"); + result->ClearUniformLoc[1] = glGetUniformLocation(result->ClearShaderPlain, "uDepth"); + result->ClearUniformLoc[2] = glGetUniformLocation(result->ClearShaderPlain, "uOpaquePolyID"); + result->ClearUniformLoc[3] = glGetUniformLocation(result->ClearShaderPlain, "uFogFlag"); memset(result->RenderShader, 0, sizeof(RenderShader)); @@ -160,42 +151,35 @@ std::unique_ptr GLRenderer::New() noexcept if (!result->BuildRenderShader(RenderFlag_ShadowMask | RenderFlag_WBuffer, kRenderVS_W, kRenderFS_WSM)) return nullptr; - if (!OpenGL::BuildShaderProgram(kFinalPassVS, kFinalPassEdgeFS, result->FinalPassEdgeShader, "FinalPassEdgeShader")) + if (!OpenGL::CompileVertexFragmentProgram(result->FinalPassEdgeShader, + kFinalPassVS, kFinalPassEdgeFS, + "FinalPassEdgeShader", + {{"vPosition", 0}}, + {{"oColor", 0}})) + return nullptr; + if (!OpenGL::CompileVertexFragmentProgram(result->FinalPassFogShader, + kFinalPassVS, kFinalPassFogFS, + "FinalPassFogShader", + {{"vPosition", 0}}, + {{"oColor", 0}})) return nullptr; - if (!OpenGL::BuildShaderProgram(kFinalPassVS, kFinalPassFogFS, result->FinalPassFogShader, "FinalPassFogShader")) - return nullptr; + GLuint uni_id = glGetUniformBlockIndex(result->FinalPassEdgeShader, "uConfig"); + glUniformBlockBinding(result->FinalPassEdgeShader, uni_id, 0); - glBindAttribLocation(result->FinalPassEdgeShader[2], 0, "vPosition"); - glBindFragDataLocation(result->FinalPassEdgeShader[2], 0, "oColor"); - - if (!OpenGL::LinkShaderProgram(result->FinalPassEdgeShader)) - return nullptr; - - GLint uni_id = glGetUniformBlockIndex(result->FinalPassEdgeShader[2], "uConfig"); - glUniformBlockBinding(result->FinalPassEdgeShader[2], uni_id, 0); - - glUseProgram(result->FinalPassEdgeShader[2]); - - uni_id = glGetUniformLocation(result->FinalPassEdgeShader[2], "DepthBuffer"); + glUseProgram(result->FinalPassEdgeShader); + uni_id = glGetUniformLocation(result->FinalPassEdgeShader, "DepthBuffer"); glUniform1i(uni_id, 0); - uni_id = glGetUniformLocation(result->FinalPassEdgeShader[2], "AttrBuffer"); + uni_id = glGetUniformLocation(result->FinalPassEdgeShader, "AttrBuffer"); glUniform1i(uni_id, 1); - glBindAttribLocation(result->FinalPassFogShader[2], 0, "vPosition"); - glBindFragDataLocation(result->FinalPassFogShader[2], 0, "oColor"); + uni_id = glGetUniformBlockIndex(result->FinalPassFogShader, "uConfig"); + glUniformBlockBinding(result->FinalPassFogShader, uni_id, 0); - if (!OpenGL::LinkShaderProgram(result->FinalPassFogShader)) - return nullptr; - - uni_id = glGetUniformBlockIndex(result->FinalPassFogShader[2], "uConfig"); - glUniformBlockBinding(result->FinalPassFogShader[2], uni_id, 0); - - glUseProgram(result->FinalPassFogShader[2]); - - uni_id = glGetUniformLocation(result->FinalPassFogShader[2], "DepthBuffer"); + glUseProgram(result->FinalPassFogShader); + uni_id = glGetUniformLocation(result->FinalPassFogShader, "DepthBuffer"); glUniform1i(uni_id, 0); - uni_id = glGetUniformLocation(result->FinalPassFogShader[2], "AttrBuffer"); + uni_id = glGetUniformLocation(result->FinalPassFogShader, "AttrBuffer"); glUniform1i(uni_id, 1); @@ -320,8 +304,8 @@ GLRenderer::~GLRenderer() for (int i = 0; i < 16; i++) { - if (!RenderShader[i][2]) continue; - OpenGL::DeleteShaderProgram(RenderShader[i]); + if (!RenderShader[i]) continue; + glDeleteProgram(RenderShader[i]); } } @@ -1093,7 +1077,7 @@ void GLRenderer::RenderSceneChunk(int y, int h) // edge marking // TODO: depth/polyid values at screen edges - glUseProgram(FinalPassEdgeShader[2]); + glUseProgram(FinalPassEdgeShader); glBlendFuncSeparate(GL_SRC_ALPHA, GL_ONE_MINUS_SRC_ALPHA, GL_ZERO, GL_ONE); @@ -1104,7 +1088,7 @@ void GLRenderer::RenderSceneChunk(int y, int h) { // fog - glUseProgram(FinalPassFogShader[2]); + glUseProgram(FinalPassFogShader); if (RenderDispCnt & (1<<6)) glBlendFuncSeparate(GL_ZERO, GL_ONE, GL_CONSTANT_COLOR, GL_ONE_MINUS_SRC_ALPHA); @@ -1238,7 +1222,7 @@ void GLRenderer::RenderFrame() // TODO: check whether 'clear polygon ID' affects translucent polyID // (for example when alpha is 1..30) { - glUseProgram(ClearShaderPlain[2]); + glUseProgram(ClearShaderPlain); glDepthFunc(GL_ALWAYS); u32 r = RenderClearAttr1 & 0x1F; diff --git a/src/GPU3D_OpenGL.h b/src/GPU3D_OpenGL.h index 597f13e1..4316b03d 100644 --- a/src/GPU3D_OpenGL.h +++ b/src/GPU3D_OpenGL.h @@ -84,13 +84,13 @@ private: }; - GLuint ClearShaderPlain[3] {}; + GLuint ClearShaderPlain {}; - GLuint RenderShader[16][3] {}; + GLuint RenderShader[16] {}; GLuint CurShaderID = -1; - GLuint FinalPassEdgeShader[3] {}; - GLuint FinalPassFogShader[3] {}; + GLuint FinalPassEdgeShader {}; + GLuint FinalPassFogShader {}; // std140 compliant structure struct diff --git a/src/GPU_OpenGL.cpp b/src/GPU_OpenGL.cpp index 47e04d25..17a5a534 100644 --- a/src/GPU_OpenGL.cpp +++ b/src/GPU_OpenGL.cpp @@ -36,32 +36,27 @@ using namespace OpenGL; std::unique_ptr GLCompositor::New() noexcept { assert(glBindAttribLocation != nullptr); + GLuint CompShader {}; - std::array CompShader {}; - if (!OpenGL::BuildShaderProgram(kCompositorVS, kCompositorFS_Nearest, &CompShader[0], "CompositorShader")) - return nullptr; - - glBindAttribLocation(CompShader[2], 0, "vPosition"); - glBindAttribLocation(CompShader[2], 1, "vTexcoord"); - glBindFragDataLocation(CompShader[2], 0, "oColor"); - - if (!OpenGL::LinkShaderProgram(CompShader.data())) - // OpenGL::LinkShaderProgram already deletes the shader program object - // if linking the shaders together failed. + if (!OpenGL::CompileVertexFragmentProgram(CompShader, + kCompositorVS, kCompositorFS_Nearest, + "CompositorShader", + {{"vPosition", 0}, {"vTexcoord", 1}}, + {{"oColor", 0}})) return nullptr; return std::unique_ptr(new GLCompositor(CompShader)); } -GLCompositor::GLCompositor(std::array compShader) noexcept : CompShader(compShader) +GLCompositor::GLCompositor(GLuint compShader) noexcept : CompShader(compShader) { - CompScaleLoc = glGetUniformLocation(CompShader[2], "u3DScale"); - Comp3DXPosLoc = glGetUniformLocation(CompShader[2], "u3DXPos"); + CompScaleLoc = glGetUniformLocation(CompShader, "u3DScale"); + Comp3DXPosLoc = glGetUniformLocation(CompShader, "u3DXPos"); - glUseProgram(CompShader[2]); - GLuint screenTextureUniform = glGetUniformLocation(CompShader[2], "ScreenTex"); + glUseProgram(CompShader); + GLuint screenTextureUniform = glGetUniformLocation(CompShader, "ScreenTex"); glUniform1i(screenTextureUniform, 0); - GLuint _3dTextureUniform = glGetUniformLocation(CompShader[2], "_3DTex"); + GLuint _3dTextureUniform = glGetUniformLocation(CompShader, "_3DTex"); glUniform1i(_3dTextureUniform, 1); // all this mess is to prevent bleeding @@ -136,7 +131,7 @@ GLCompositor::~GLCompositor() glDeleteVertexArrays(1, &CompVertexArrayID); glDeleteBuffers(1, &CompVertexBufferID); - OpenGL::DeleteShaderProgram(CompShader.data()); + glDeleteProgram(CompShader); } void GLCompositor::Reset() @@ -200,7 +195,7 @@ void GLCompositor::RenderFrame() glClear(GL_COLOR_BUFFER_BIT); // TODO: select more shaders (filtering, etc) - OpenGL::UseShaderProgram(CompShader.data()); + glUseProgram(CompShader); glUniform1ui(CompScaleLoc, Scale); // TODO: support setting this midframe, if ever needed diff --git a/src/GPU_OpenGL.h b/src/GPU_OpenGL.h index 90c17ae3..bf4e9d3d 100644 --- a/src/GPU_OpenGL.h +++ b/src/GPU_OpenGL.h @@ -44,12 +44,12 @@ public: void RenderFrame(); void BindOutputTexture(int buf); private: - GLCompositor(std::array CompShader) noexcept; + GLCompositor(GLuint CompShader) noexcept; int Scale; int ScreenH, ScreenW; - std::array CompShader; + GLuint CompShader; GLuint CompScaleLoc; GLuint Comp3DXPosLoc; diff --git a/src/NonStupidBitfield.h b/src/NonStupidBitfield.h index a4fe7ec8..ad76c64d 100644 --- a/src/NonStupidBitfield.h +++ b/src/NonStupidBitfield.h @@ -26,6 +26,32 @@ #include #include +inline u64 GetRangedBitMask(u32 idx, u32 startBit, u32 bitsCount) +{ + u32 startEntry = startBit >> 6; + u64 entriesCount = ((startBit + bitsCount + 0x3F) >> 6) - startEntry; + + if (entriesCount > 1) + { + if (idx == startEntry) + return 0xFFFFFFFFFFFFFFFF << (startBit & 0x3F); + if (((startBit + bitsCount) & 0x3F) && idx == startEntry + entriesCount - 1) + return ~(0xFFFFFFFFFFFFFFFF << ((startBit + bitsCount) & 0x3F)); + else + return 0xFFFFFFFFFFFFFFFF; + + return 0xFFFFFFFFFFFFFFFF; + } + else if (idx == startEntry) + { + return ((1ULL << bitsCount) - 1) << (startBit & 0x3F); + } + else + { + return 0; + } +} + // like std::bitset but less stupid and optimised for // our use case (keeping track of memory invalidations) @@ -164,6 +190,11 @@ struct NonStupidBitField return Ref{*this, idx}; } + bool operator[](u32 idx) const + { + return Data[idx >> 6] & (1ULL << (idx & 0x3F)); + } + void SetRange(u32 startBit, u32 bitsCount) { u32 startEntry = startBit >> 6; @@ -185,6 +216,26 @@ struct NonStupidBitField } } + int Min() const + { + for (int i = 0; i < DataLength; i++) + { + if (Data[i]) + return i * 64 + __builtin_ctzll(Data[i]); + } + return -1; + } + + int Max() const + { + for (int i = DataLength - 1; i >= 0; i--) + { + if (Data[i]) + return i * 64 + (63 - __builtin_clzll(Data[i])); + } + return -1; + } + NonStupidBitField& operator|=(const NonStupidBitField& other) { for (u32 i = 0; i < DataLength; i++) @@ -193,6 +244,7 @@ struct NonStupidBitField } return *this; } + NonStupidBitField& operator&=(const NonStupidBitField& other) { for (u32 i = 0; i < DataLength; i++) @@ -201,6 +253,20 @@ struct NonStupidBitField } return *this; } + + operator bool() const + { + for (int i = 0; i < DataLength - 1; i++) + { + if (Data[i]) + return true; + } + if (Data[DataLength-1] & ((Size&0x3F) ? ~(0xFFFFFFFFFFFFFFFF << (Size&0x3F)) : 0xFFFFFFFFFFFFFFFF)) + { + return true; + } + return false; + } }; diff --git a/src/OpenGLSupport.cpp b/src/OpenGLSupport.cpp index f1914fc1..6f6d770b 100644 --- a/src/OpenGLSupport.cpp +++ b/src/OpenGLSupport.cpp @@ -24,7 +24,9 @@ using Platform::LogLevel; namespace OpenGL { -bool BuildShaderProgram(const char* vs, const char* fs, GLuint* ids, const char* name) +#define checkGLError() if (glGetError() != GL_NO_ERROR) printf("error %d\n", __LINE__) + +bool CompilerShader(GLuint& id, const char* source, const char* name, const char* type) { int len; int res; @@ -35,61 +37,31 @@ bool BuildShaderProgram(const char* vs, const char* fs, GLuint* ids, const char* return false; } - ids[0] = glCreateShader(GL_VERTEX_SHADER); - len = strlen(vs); - glShaderSource(ids[0], 1, &vs, &len); - glCompileShader(ids[0]); + len = strlen(source); + glShaderSource(id, 1, &source, &len); + checkGLError(); + glCompileShader(id); + checkGLError(); - glGetShaderiv(ids[0], GL_COMPILE_STATUS, &res); + glGetShaderiv(id, GL_COMPILE_STATUS, &res); + checkGLError(); if (res != GL_TRUE) { - glGetShaderiv(ids[0], GL_INFO_LOG_LENGTH, &res); + glGetShaderiv(id, GL_INFO_LOG_LENGTH, &res); if (res < 1) res = 1024; char* log = new char[res+1]; - glGetShaderInfoLog(ids[0], res+1, NULL, log); - Log(LogLevel::Error, "OpenGL: failed to compile vertex shader %s: %s\n", name, log); - Log(LogLevel::Debug, "shader source:\n--\n%s\n--\n", vs); + glGetShaderInfoLog(id, res+1, NULL, log); + Log(LogLevel::Error, "OpenGL: failed to compile %s shader %s: %s\n", type, name, log); + Log(LogLevel::Debug, "shader source:\n--\n%s\n--\n", source); delete[] log; - glDeleteShader(ids[0]); - return false; } - ids[1] = glCreateShader(GL_FRAGMENT_SHADER); - len = strlen(fs); - glShaderSource(ids[1], 1, &fs, &len); - glCompileShader(ids[1]); - - glGetShaderiv(ids[1], GL_COMPILE_STATUS, &res); - if (res != GL_TRUE) - { - glGetShaderiv(ids[1], GL_INFO_LOG_LENGTH, &res); - if (res < 1) res = 1024; - char* log = new char[res+1]; - glGetShaderInfoLog(ids[1], res+1, NULL, log); - Log(LogLevel::Error, "OpenGL: failed to compile fragment shader %s: %s\n", name, log); - //printf("shader source:\n--\n%s\n--\n", fs); - delete[] log; - - FILE* logf = fopen("shaderfail.log", "w"); - fwrite(fs, len+1, 1, logf); - fclose(logf); - - glDeleteShader(ids[0]); - glDeleteShader(ids[1]); - - return false; - } - - ids[2] = glCreateProgram(); - glAttachShader(ids[2], ids[0]); - glAttachShader(ids[2], ids[1]); - return true; } -bool LinkShaderProgram(GLuint* ids) +bool LinkProgram(GLuint& result, GLuint* ids, int numIds) { int res; @@ -99,46 +71,100 @@ bool LinkShaderProgram(GLuint* ids) return false; } - glLinkProgram(ids[2]); + for (int i = 0; i < numIds; i++) + { + glAttachShader(result, ids[i]); + checkGLError(); + } - glDetachShader(ids[2], ids[0]); - glDetachShader(ids[2], ids[1]); + glLinkProgram(result); - glDeleteShader(ids[0]); - glDeleteShader(ids[1]); + for (int i = 0; i < numIds; i++) + glDetachShader(result, ids[i]); - glGetProgramiv(ids[2], GL_LINK_STATUS, &res); + glGetProgramiv(result, GL_LINK_STATUS, &res); if (res != GL_TRUE) { - glGetProgramiv(ids[2], GL_INFO_LOG_LENGTH, &res); + glGetProgramiv(result, GL_INFO_LOG_LENGTH, &res); if (res < 1) res = 1024; char* log = new char[res+1]; - glGetProgramInfoLog(ids[2], res+1, NULL, log); + glGetProgramInfoLog(result, res+1, NULL, log); Log(LogLevel::Error, "OpenGL: failed to link shader program: %s\n", log); delete[] log; - glDeleteProgram(ids[2]); - return false; } return true; } -void DeleteShaderProgram(GLuint* ids) +bool CompileComputeProgram(GLuint& result, const char* source, const char* name) { + GLuint shader = glCreateShader(GL_COMPUTE_SHADER); + bool linkingSucess = false; if (glDeleteProgram) { // If OpenGL isn't loaded, then there's no shader program to delete - glDeleteProgram(ids[2]); + goto error; } + + result = glCreateProgram(); + + printf("compiling %s", name); + if (!CompilerShader(shader, source, name, "compute")) + goto error; + + linkingSucess = LinkProgram(result, &shader, 1); + +error: + glDeleteShader(shader); + + if (!linkingSucess) + glDeleteProgram(result); + + return linkingSucess; } -void UseShaderProgram(GLuint* ids) +bool CompileVertexFragmentProgram(GLuint& result, + const char* vs, const char* fs, + const char* name, + const std::initializer_list& vertexInAttrs, + const std::initializer_list& fragmentOutAttrs) { - if (glUseProgram) - { // If OpenGL isn't loaded, then there's no shader program to use - glUseProgram(ids[2]); + GLuint shaders[2] = + { + glCreateShader(GL_VERTEX_SHADER), + glCreateShader(GL_FRAGMENT_SHADER) + }; + result = glCreateProgram(); + + bool linkingSucess = false; + + if (!CompilerShader(shaders[0], vs, name, "vertex")) + goto error; + + if (!CompilerShader(shaders[1], fs, name, "fragment")) + goto error; + + + for (const AttributeTarget& target : vertexInAttrs) + { + glBindAttribLocation(result, target.Location, target.Name); } + for (const AttributeTarget& target : fragmentOutAttrs) + { + glBindFragDataLocation(result, target.Location, target.Name); + } + + linkingSucess = LinkProgram(result, shaders, 2); + +error: + glDeleteShader(shaders[1]); + glDeleteShader(shaders[0]); + + if (!linkingSucess) + glDeleteProgram(result); + + return linkingSucess; } } diff --git a/src/OpenGLSupport.h b/src/OpenGLSupport.h index 14be01a6..7d8aae44 100644 --- a/src/OpenGLSupport.h +++ b/src/OpenGLSupport.h @@ -29,10 +29,19 @@ namespace OpenGL { -bool BuildShaderProgram(const char* vs, const char* fs, GLuint* ids, const char* name); -bool LinkShaderProgram(GLuint* ids); -void DeleteShaderProgram(GLuint* ids); -void UseShaderProgram(GLuint* ids); +struct AttributeTarget +{ + const char* Name; + u32 Location; +}; + +bool CompileVertexFragmentProgram(GLuint& result, + const char* vs, const char* fs, + const char* name, + const std::initializer_list& vertexInAttrs, + const std::initializer_list& fragmentOutAttrs); + +bool CompileComputeProgram(GLuint& result, const char* source, const char* name); } diff --git a/src/frontend/qt_sdl/CMakeLists.txt b/src/frontend/qt_sdl/CMakeLists.txt index 24261030..a5f68033 100644 --- a/src/frontend/qt_sdl/CMakeLists.txt +++ b/src/frontend/qt_sdl/CMakeLists.txt @@ -24,6 +24,7 @@ set(SOURCES_QT_SDL ROMInfoDialog.cpp RAMInfoDialog.cpp TitleManagerDialog.cpp + RTCSettingsDialog.cpp Input.cpp LAN_PCap.cpp LAN_Socket.cpp diff --git a/src/frontend/qt_sdl/OSD.cpp b/src/frontend/qt_sdl/OSD.cpp index d3becc12..6842d5f0 100644 --- a/src/frontend/qt_sdl/OSD.cpp +++ b/src/frontend/qt_sdl/OSD.cpp @@ -57,7 +57,7 @@ struct Item std::deque ItemQueue; -GLuint Shader[3]; +GLuint Shader; GLint uScreenSize, uOSDPos, uOSDSize; GLfloat uScaleFactor; GLuint OSDVertexArray; @@ -70,20 +70,19 @@ bool Init(bool openGL) { if (openGL) { - OpenGL::BuildShaderProgram(kScreenVS_OSD, kScreenFS_OSD, Shader, "OSDShader"); + OpenGL::CompileVertexFragmentProgram(Shader, + kScreenVS_OSD, kScreenFS_OSD, + "OSDShader", + {{"vPosition", 0}}, + {{"oColor", 0}}); - GLuint pid = Shader[2]; - glBindAttribLocation(pid, 0, "vPosition"); - glBindFragDataLocation(pid, 0, "oColor"); + glUseProgram(Shader); + glUniform1i(glGetUniformLocation(Shader, "OSDTex"), 0); - OpenGL::LinkShaderProgram(Shader); - glUseProgram(pid); - glUniform1i(glGetUniformLocation(pid, "OSDTex"), 0); - - uScreenSize = glGetUniformLocation(pid, "uScreenSize"); - uOSDPos = glGetUniformLocation(pid, "uOSDPos"); - uOSDSize = glGetUniformLocation(pid, "uOSDSize"); - uScaleFactor = glGetUniformLocation(pid, "uScaleFactor"); + uScreenSize = glGetUniformLocation(Shader, "uScreenSize"); + uOSDPos = glGetUniformLocation(Shader, "uOSDPos"); + uOSDSize = glGetUniformLocation(Shader, "uOSDSize"); + uScaleFactor = glGetUniformLocation(Shader, "uScaleFactor"); float vertices[6*2] = { @@ -425,7 +424,7 @@ void DrawGL(float w, float h) u32 y = kOSDMargin; - glUseProgram(Shader[2]); + glUseProgram(Shader); glUniform2f(uScreenSize, w, h); glUniform1f(uScaleFactor, mainWindow->devicePixelRatioF()); diff --git a/src/frontend/qt_sdl/main.cpp b/src/frontend/qt_sdl/main.cpp index a6cc2482..875a39b6 100644 --- a/src/frontend/qt_sdl/main.cpp +++ b/src/frontend/qt_sdl/main.cpp @@ -228,19 +228,17 @@ void EmuThread::initOpenGL() oglContext = windowctx; oglContext->MakeCurrent(); - OpenGL::BuildShaderProgram(kScreenVS, kScreenFS, screenShaderProgram, "ScreenShader"); - GLuint pid = screenShaderProgram[2]; - glBindAttribLocation(pid, 0, "vPosition"); - glBindAttribLocation(pid, 1, "vTexcoord"); - glBindFragDataLocation(pid, 0, "oColor"); + OpenGL::CompileVertexFragmentProgram(screenShaderProgram, + kScreenVS, kScreenFS, + "ScreenShader", + {{"vPosition", 0}, {"vTexcoord", 1}}, + {{"oColor", 0}}); - OpenGL::LinkShaderProgram(screenShaderProgram); + glUseProgram(screenShaderProgram); + glUniform1i(glGetUniformLocation(screenShaderProgram, "ScreenTex"), 0); - glUseProgram(pid); - glUniform1i(glGetUniformLocation(pid, "ScreenTex"), 0); - - screenShaderScreenSizeULoc = glGetUniformLocation(pid, "uScreenSize"); - screenShaderTransformULoc = glGetUniformLocation(pid, "uTransform"); + screenShaderScreenSizeULoc = glGetUniformLocation(screenShaderProgram, "uScreenSize"); + screenShaderTransformULoc = glGetUniformLocation(screenShaderProgram, "uTransform"); // to prevent bleeding between both parts of the screen // with bilinear filtering enabled @@ -300,7 +298,7 @@ void EmuThread::deinitOpenGL() glDeleteVertexArrays(1, &screenVertexArray); glDeleteBuffers(1, &screenVertexBuffer); - OpenGL::DeleteShaderProgram(screenShaderProgram); + glDeleteProgram(screenShaderProgram); OSD::DeInit(); @@ -749,7 +747,7 @@ void EmuThread::drawScreenGL() glViewport(0, 0, w, h); - glUseProgram(screenShaderProgram[2]); + glUseProgram(screenShaderProgram); glUniform2f(screenShaderScreenSizeULoc, w / factor, h / factor); int frontbuf = FrontBuffer; diff --git a/src/frontend/qt_sdl/main.h b/src/frontend/qt_sdl/main.h index 073a4da0..aa94a625 100644 --- a/src/frontend/qt_sdl/main.h +++ b/src/frontend/qt_sdl/main.h @@ -121,7 +121,7 @@ private: GL::Context* oglContext = nullptr; GLuint screenVertexBuffer, screenVertexArray; GLuint screenTexture; - GLuint screenShaderProgram[3]; + GLuint screenShaderProgram; GLuint screenShaderTransformULoc, screenShaderScreenSizeULoc; QMutex screenSettingsLock;