Texture Cache: Improve documentation

This commit adds a series of HLE methods for handling 3D textures in
general. This helps games that generate 3D textures on every frame and
may reduce loading times for certain games.

.appveyor/UtilityFunctions.ps1

@@ -1,39 +0,0 @@
-# Set-up Visual Studio Command Prompt environment for PowerShell
-pushd "C:\Program Files (x86)\Microsoft Visual Studio\2017\Community\Common7\Tools\"
-cmd /c "VsDevCmd.bat -arch=x64 & set" | foreach {
-    if ($_ -match "=") {
-        $v = $_.split("="); Set-Item -Force -Path "ENV:\$($v[0])" -Value "$($v[1])"
-    }
-}
-popd
-
-function Which ($search_path, $name) {
-    ($search_path).Split(";") | Get-ChildItem -Filter $name | Select -First 1 -Exp FullName
-}
-
-function GetDeps ($search_path, $binary) {
-    ((dumpbin /dependents $binary).Where({ $_ -match "dependencies:"}, "SkipUntil") | Select-String "[^ ]*\.dll").Matches | foreach {
-        Which $search_path $_.Value
-    }
-}
-
-function RecursivelyGetDeps ($search_path, $binary) {
-    $final_deps = @()
-    $deps_to_process = GetDeps $search_path $binary
-    while ($deps_to_process.Count -gt 0) {
-        $current, $deps_to_process = $deps_to_process
-        if ($final_deps -contains $current) { continue }
-
-        # Is this a system dll file?
-        # We use the same algorithm that cmake uses to determine this.
-        if ($current -match "$([regex]::Escape($env:SystemRoot))\\sys") { continue }
-        if ($current -match "$([regex]::Escape($env:WinDir))\\sys") { continue }
-        if ($current -match "\\msvc[^\\]+dll") { continue }
-        if ($current -match "\\api-ms-win-[^\\]+dll") { continue }
-
-        $final_deps += $current
-        $new_deps = GetDeps $search_path $current
-        $deps_to_process += ($new_deps | ?{-not ($final_deps -contains $_)})
-    }
-    return $final_deps
-}

.ci/scripts/windows/upload.ps1

@@ -39,6 +39,7 @@ mkdir "artifacts"
 # Build a tar.xz for the source of the release
 Copy-Item .\license.txt -Destination $MSVC_SOURCE
 Copy-Item .\README.md -Destination $MSVC_SOURCE
+Copy-Item .\CMakeLists.txt -Destination $MSVC_SOURCE
 Copy-Item .\src -Recurse -Destination $MSVC_SOURCE
 Copy-Item .\externals -Recurse -Destination $MSVC_SOURCE
 Copy-Item .\dist -Recurse -Destination $MSVC_SOURCE
@@ -60,4 +61,4 @@ Get-ChildItem "$BUILD_DIR" -Recurse -Filter "QtWebEngineProcess*.exe" | Copy-Ite
 
 Get-ChildItem . -Filter "*.zip" | Copy-Item -destination "artifacts"
 Get-ChildItem . -Filter "*.7z" | Copy-Item -destination "artifacts"
-Get-ChildItem . -Filter "*.tar.xz" | Copy-Item -destination "artifacts"
+Get-ChildItem . -Filter "*.tar.xz" | Copy-Item -destination "artifacts"

CMakeLists.txt

@@ -29,7 +29,7 @@ option(ENABLE_VULKAN "Enables Vulkan backend" ON)
 
 option(USE_DISCORD_PRESENCE "Enables Discord Rich Presence" OFF)
 
-if(NOT EXISTS ${PROJECT_SOURCE_DIR}/.git/hooks/pre-commit)
+if(EXISTS ${PROJECT_SOURCE_DIR}/hooks/pre-commit AND NOT EXISTS ${PROJECT_SOURCE_DIR}/.git/hooks/pre-commit)
     message(STATUS "Copying pre-commit hook")
     file(COPY hooks/pre-commit
         DESTINATION ${PROJECT_SOURCE_DIR}/.git/hooks)
@@ -49,7 +49,10 @@ function(check_submodules_present)
         endif()
     endforeach()
 endfunction()
-check_submodules_present()
+
+if(EXISTS ${PROJECT_SOURCE_DIR}/.gitmodules)
+    check_submodules_present()
+endif()
 
 configure_file(${PROJECT_SOURCE_DIR}/dist/compatibility_list/compatibility_list.qrc
                ${PROJECT_BINARY_DIR}/dist/compatibility_list/compatibility_list.qrc

appveyor.yml

@@ -1,178 +0,0 @@
-# shallow clone
-clone_depth: 10
-
-cache:
-  - C:\ProgramData\chocolatey\bin -> appveyor.yml
-  - C:\ProgramData\chocolatey\lib -> appveyor.yml
-
-os: Visual Studio 2017
-
-environment:
-  # Tell msys2 to add mingw64 to the path
-  MSYSTEM: MINGW64
-  # Tell msys2 to inherit the current directory when starting the shell
-  CHERE_INVOKING: 1
-  matrix:
-    - BUILD_TYPE: msvc
-    - BUILD_TYPE: mingw
-
-platform:
-  - x64
-
-configuration:
-  - Release
-
-install:
-  - git submodule update --init --recursive
-  - ps: |
-        if ($env:BUILD_TYPE -eq 'mingw') {
-          $dependencies = "mingw64/mingw-w64-x86_64-cmake",
-                          "mingw64/mingw-w64-x86_64-qt5",
-                          "mingw64/mingw-w64-x86_64-SDL2"
-          # redirect err to null to prevent warnings from becoming errors
-          # workaround to prevent pacman from failing due to cyclical dependencies
-          C:\msys64\usr\bin\bash -lc "pacman --noconfirm -S mingw64/mingw-w64-x86_64-freetype mingw64/mingw-w64-x86_64-fontconfig" 2> $null
-          C:\msys64\usr\bin\bash -lc "pacman --noconfirm -S $dependencies" 2> $null
-        }
-
-before_build:
-  - mkdir %BUILD_TYPE%_build
-  - cd %BUILD_TYPE%_build
-  - ps: |
-        $COMPAT = if ($env:ENABLE_COMPATIBILITY_REPORTING -eq $null) {0} else {$env:ENABLE_COMPATIBILITY_REPORTING}
-        if ($env:BUILD_TYPE -eq 'msvc') {
-          # redirect stderr and change the exit code to prevent powershell from cancelling the build if cmake prints a warning
-          cmd /C 'cmake -G "Visual Studio 15 2017 Win64" -DYUZU_USE_BUNDLED_QT=1 -DYUZU_USE_BUNDLED_SDL2=1 -DYUZU_USE_BUNDLED_UNICORN=1 -DYUZU_USE_QT_WEB_ENGINE=ON -DENABLE_COMPATIBILITY_LIST_DOWNLOAD=ON -DYUZU_ENABLE_COMPATIBILITY_REPORTING=${COMPAT} -DUSE_DISCORD_PRESENCE=ON .. 2>&1 && exit 0'
-        } else {
-          C:\msys64\usr\bin\bash.exe -lc "cmake -G 'MSYS Makefiles' -DYUZU_BUILD_UNICORN=1 -DCMAKE_BUILD_TYPE=Release -DENABLE_COMPATIBILITY_LIST_DOWNLOAD=ON -DYUZU_ENABLE_COMPATIBILITY_REPORTING=${COMPAT} -DUSE_DISCORD_PRESENCE=ON .. 2>&1"
-        }
-  - cd ..
-
-build_script:
-  - ps: |
-        if ($env:BUILD_TYPE -eq 'msvc') {
-          # https://www.appveyor.com/docs/build-phase
-          msbuild msvc_build/yuzu.sln /maxcpucount /logger:"C:\Program Files\AppVeyor\BuildAgent\Appveyor.MSBuildLogger.dll"
-        } else {
-          C:\msys64\usr\bin\bash.exe -lc 'mingw32-make -C mingw_build/ 2>&1'
-        }
-
-after_build:
-  - ps: |
-        $GITDATE = $(git show -s --date=short --format='%ad') -replace "-",""
-        $GITREV = $(git show -s --format='%h')
-
-        # Find out which kind of release we are producing by tag name
-        if ($env:APPVEYOR_REPO_TAG_NAME) {
-          $RELEASE_DIST, $RELEASE_VERSION = $env:APPVEYOR_REPO_TAG_NAME.split('-')
-        } else {
-          # There is no repo tag - make assumptions
-          $RELEASE_DIST = "head"
-        }
-
-        if ($env:BUILD_TYPE -eq 'msvc') {
-          # Where are these spaces coming from? Regardless, let's remove them
-          $MSVC_BUILD_ZIP = "yuzu-windows-msvc-$GITDATE-$GITREV.zip" -replace " ", ""
-          $MSVC_BUILD_PDB = "yuzu-windows-msvc-$GITDATE-$GITREV-debugsymbols.zip" -replace " ", ""
-          $MSVC_SEVENZIP = "yuzu-windows-msvc-$GITDATE-$GITREV.7z" -replace " ", ""
-
-          # set the build names as env vars so the artifacts can upload them
-          $env:BUILD_ZIP = $MSVC_BUILD_ZIP
-          $env:BUILD_SYMBOLS = $MSVC_BUILD_PDB
-          $env:BUILD_UPDATE = $MSVC_SEVENZIP
-
-          $BUILD_DIR = ".\msvc_build\bin\Release"
-
-          # Make a debug symbol upload
-          mkdir pdb
-          Get-ChildItem "$BUILD_DIR\" -Recurse -Filter "*.pdb" | Copy-Item -destination .\pdb
-          7z a -tzip $MSVC_BUILD_PDB .\pdb\*.pdb
-          rm "$BUILD_DIR\*.pdb"
-
-          mkdir $RELEASE_DIST
-          # get rid of extra exes by copying everything over, then deleting all the exes, then copying just the exes we want
-          Copy-Item "$BUILD_DIR\*" -Destination $RELEASE_DIST -Recurse
-          rm "$RELEASE_DIST\*.exe"
-          Get-ChildItem "$BUILD_DIR" -Recurse -Filter "yuzu*.exe" | Copy-Item -destination $RELEASE_DIST
-          Get-ChildItem "$BUILD_DIR" -Recurse -Filter "QtWebEngineProcess*.exe" | Copy-Item -destination $RELEASE_DIST
-          Copy-Item .\license.txt -Destination $RELEASE_DIST
-          Copy-Item .\README.md -Destination $RELEASE_DIST
-          7z a -tzip $MSVC_BUILD_ZIP $RELEASE_DIST\*
-          7z a $MSVC_SEVENZIP $RELEASE_DIST
-        } else {
-          $MINGW_BUILD_ZIP = "yuzu-windows-mingw-$GITDATE-$GITREV.zip" -replace " ", ""
-          $MINGW_SEVENZIP = "yuzu-windows-mingw-$GITDATE-$GITREV.7z" -replace " ", ""
-          # not going to bother adding separate debug symbols for mingw, so just upload a README for it
-          # if someone wants to add them, change mingw to compile with -g and use objdump and strip to separate the symbols from the binary
-          $MINGW_NO_DEBUG_SYMBOLS = "README_No_Debug_Symbols.txt"
-          Set-Content -Path $MINGW_NO_DEBUG_SYMBOLS -Value "This is a workaround for Appveyor since msvc has debug symbols but mingw doesnt" -Force
-
-          # store the build information in env vars so we can use them as artifacts
-          $env:BUILD_ZIP = $MINGW_BUILD_ZIP
-          $env:BUILD_SYMBOLS = $MINGW_NO_DEBUG_SYMBOLS
-          $env:BUILD_UPDATE = $MINGW_SEVENZIP
-
-          $CMAKE_SOURCE_DIR = "$env:APPVEYOR_BUILD_FOLDER"
-          $CMAKE_BINARY_DIR = "$CMAKE_SOURCE_DIR/mingw_build/bin"
-          $RELEASE_DIST = $RELEASE_DIST + "-mingw"
-
-          mkdir $RELEASE_DIST
-          mkdir $RELEASE_DIST/platforms
-          mkdir $RELEASE_DIST/styles
-          mkdir $RELEASE_DIST/imageformats
-
-          # copy the compiled binaries and other release files to the release folder
-          Get-ChildItem "$CMAKE_BINARY_DIR" -Filter "yuzu*.exe" | Copy-Item -destination $RELEASE_DIST
-          Copy-Item -path "$CMAKE_SOURCE_DIR/license.txt" -destination $RELEASE_DIST
-          Copy-Item -path "$CMAKE_SOURCE_DIR/README.md" -destination $RELEASE_DIST
-
-          # copy the qt windows plugin dll to platforms
-          Copy-Item -path "C:/msys64/mingw64/share/qt5/plugins/platforms/qwindows.dll" -force -destination "$RELEASE_DIST/platforms"
-
-          # copy the qt windows vista style dll to platforms
-          Copy-Item -path "C:/msys64/mingw64/share/qt5/plugins/styles/qwindowsvistastyle.dll" -force -destination "$RELEASE_DIST/styles"
-
-          # copy the qt jpeg imageformat dll to platforms
-          Copy-Item -path "C:/msys64/mingw64/share/qt5/plugins/imageformats/qjpeg.dll" -force -destination "$RELEASE_DIST/imageformats"
-
-          # copy all the dll dependencies to the release folder
-          . "./.appveyor/UtilityFunctions.ps1"
-          $DLLSearchPath = "C:\msys64\mingw64\bin;$env:PATH"
-          $MingwDLLs = RecursivelyGetDeps $DLLSearchPath "$RELEASE_DIST\yuzu.exe"
-          $MingwDLLs += RecursivelyGetDeps $DLLSearchPath  "$RELEASE_DIST\yuzu_cmd.exe"
-          $MingwDLLs += RecursivelyGetDeps $DLLSearchPath  "$RELEASE_DIST\imageformats\qjpeg.dll"
-          Write-Host "Detected the following dependencies:"
-          Write-Host $MingwDLLs
-          foreach ($file in $MingwDLLs) {
-            Copy-Item -path "$file" -force -destination "$RELEASE_DIST"
-          }
-
-          7z a -tzip $MINGW_BUILD_ZIP $RELEASE_DIST\*
-          7z a $MINGW_SEVENZIP $RELEASE_DIST
-        }
-
-test_script:
-  - cd %BUILD_TYPE%_build
-  - ps: |
-        if ($env:BUILD_TYPE -eq 'msvc') {
-          ctest -VV -C Release
-        } else {
-          C:\msys64\usr\bin\bash.exe -lc "ctest -VV -C Release"
-        }
-  - cd ..
-
-artifacts:
-  - path: $(BUILD_ZIP)
-    name: build
-    type: zip
-
-deploy:
-  provider: GitHub
-  release: $(appveyor_repo_tag_name)
-  auth_token:
-    secure: QqePPnXbkzmXct5c8hZ2X5AbsthbI6cS1Sr+VBzcD8oUOIjfWJJKXVAQGUbQAbb0
-  artifact: update,build
-  draft: false
-  prerelease: false
-  on:
-    appveyor_repo_tag: true

src/common/threadsafe_queue.h

@@ -46,9 +46,16 @@ public:
         ElementPtr* new_ptr = new ElementPtr();
         write_ptr->next.store(new_ptr, std::memory_order_release);
         write_ptr = new_ptr;
-        cv.notify_one();
 
-        ++size;
+        const size_t previous_size{size++};
+
+        // Acquire the mutex and then immediately release it as a fence.
+        // TODO(bunnei): This can be replaced with C++20 waitable atomics when properly supported.
+        // See discussion on https://github.com/yuzu-emu/yuzu/pull/3173 for details.
+        if (previous_size == 0) {
+            std::lock_guard lock{cv_mutex};
+        }
+        cv.notify_one();
     }
 
     void Pop() {

src/core/CMakeLists.txt

@@ -170,6 +170,7 @@ add_library(core STATIC
     hle/kernel/server_port.h
     hle/kernel/server_session.cpp
     hle/kernel/server_session.h
+    hle/kernel/session.cpp
     hle/kernel/session.h
     hle/kernel/shared_memory.cpp
     hle/kernel/shared_memory.h

src/core/core_cpu.cpp

@@ -96,6 +96,8 @@ void Cpu::RunLoop(bool tight_loop) {
         } else {
             arm_interface->Step();
         }
+        // We are stopping a run, exclusive state must be cleared
+        arm_interface->ClearExclusiveState();
     }
     core_timing.Advance();
 

src/core/hle/ipc_helpers.h

@@ -19,6 +19,7 @@
 #include "core/hle/kernel/hle_ipc.h"
 #include "core/hle/kernel/object.h"
 #include "core/hle/kernel/server_session.h"
+#include "core/hle/kernel/session.h"
 #include "core/hle/result.h"
 
 namespace IPC {
@@ -139,10 +140,9 @@ public:
             context->AddDomainObject(std::move(iface));
         } else {
             auto& kernel = Core::System::GetInstance().Kernel();
-            auto [server, client] =
-                Kernel::ServerSession::CreateSessionPair(kernel, iface->GetServiceName());
-            iface->ClientConnected(server);
+            auto [client, server] = Kernel::Session::Create(kernel, iface->GetServiceName());
             context->AddMoveObject(std::move(client));
+            iface->ClientConnected(std::move(server));
         }
     }
 

src/core/hle/kernel/address_arbiter.cpp

@@ -11,18 +11,16 @@
 #include "core/core_cpu.h"
 #include "core/hle/kernel/address_arbiter.h"
 #include "core/hle/kernel/errors.h"
-#include "core/hle/kernel/object.h"
-#include "core/hle/kernel/process.h"
 #include "core/hle/kernel/scheduler.h"
 #include "core/hle/kernel/thread.h"
 #include "core/hle/result.h"
 #include "core/memory.h"
 
 namespace Kernel {
-namespace {
+
 // Wake up num_to_wake (or all) threads in a vector.
-void WakeThreads(const std::vector<std::shared_ptr<Thread>>& waiting_threads, s32 num_to_wake) {
-    auto& system = Core::System::GetInstance();
+void AddressArbiter::WakeThreads(const std::vector<std::shared_ptr<Thread>>& waiting_threads,
+                                 s32 num_to_wake) {
     // Only process up to 'target' threads, unless 'target' is <= 0, in which case process
     // them all.
     std::size_t last = waiting_threads.size();
@@ -34,12 +32,12 @@ void WakeThreads(const std::vector<std::shared_ptr<Thread>>& waiting_threads, s3
     for (std::size_t i = 0; i < last; i++) {
         ASSERT(waiting_threads[i]->GetStatus() == ThreadStatus::WaitArb);
         waiting_threads[i]->SetWaitSynchronizationResult(RESULT_SUCCESS);
+        RemoveThread(waiting_threads[i]);
         waiting_threads[i]->SetArbiterWaitAddress(0);
         waiting_threads[i]->ResumeFromWait();
         system.PrepareReschedule(waiting_threads[i]->GetProcessorID());
     }
 }
-} // Anonymous namespace
 
 AddressArbiter::AddressArbiter(Core::System& system) : system{system} {}
 AddressArbiter::~AddressArbiter() = default;
@@ -186,6 +184,7 @@ ResultCode AddressArbiter::WaitForAddressIfEqual(VAddr address, s32 value, s64 t
 ResultCode AddressArbiter::WaitForAddressImpl(VAddr address, s64 timeout) {
     Thread* current_thread = system.CurrentScheduler().GetCurrentThread();
     current_thread->SetArbiterWaitAddress(address);
+    InsertThread(SharedFrom(current_thread));
     current_thread->SetStatus(ThreadStatus::WaitArb);
     current_thread->InvalidateWakeupCallback();
     current_thread->WakeAfterDelay(timeout);
@@ -194,26 +193,51 @@ ResultCode AddressArbiter::WaitForAddressImpl(VAddr address, s64 timeout) {
     return RESULT_TIMEOUT;
 }
 
-std::vector<std::shared_ptr<Thread>> AddressArbiter::GetThreadsWaitingOnAddress(
-    VAddr address) const {
+void AddressArbiter::HandleWakeupThread(std::shared_ptr<Thread> thread) {
+    ASSERT(thread->GetStatus() == ThreadStatus::WaitArb);
+    RemoveThread(thread);
+    thread->SetArbiterWaitAddress(0);
+}
 
-    // Retrieve all threads that are waiting for this address.
-    std::vector<std::shared_ptr<Thread>> threads;
-    const auto& scheduler = system.GlobalScheduler();
-    const auto& thread_list = scheduler.GetThreadList();
-
-    for (const auto& thread : thread_list) {
-        if (thread->GetArbiterWaitAddress() == address) {
-            threads.push_back(thread);
+void AddressArbiter::InsertThread(std::shared_ptr<Thread> thread) {
+    const VAddr arb_addr = thread->GetArbiterWaitAddress();
+    std::list<std::shared_ptr<Thread>>& thread_list = arb_threads[arb_addr];
+    auto it = thread_list.begin();
+    while (it != thread_list.end()) {
+        const std::shared_ptr<Thread>& current_thread = *it;
+        if (current_thread->GetPriority() >= thread->GetPriority()) {
+            thread_list.insert(it, thread);
+            return;
         }
+        ++it;
     }
+    thread_list.push_back(std::move(thread));
+}
 
-    // Sort them by priority, such that the highest priority ones come first.
-    std::sort(threads.begin(), threads.end(),
-              [](const std::shared_ptr<Thread>& lhs, const std::shared_ptr<Thread>& rhs) {
-                  return lhs->GetPriority() < rhs->GetPriority();
-              });
+void AddressArbiter::RemoveThread(std::shared_ptr<Thread> thread) {
+    const VAddr arb_addr = thread->GetArbiterWaitAddress();
+    std::list<std::shared_ptr<Thread>>& thread_list = arb_threads[arb_addr];
+    auto it = thread_list.begin();
+    while (it != thread_list.end()) {
+        const std::shared_ptr<Thread>& current_thread = *it;
+        if (current_thread.get() == thread.get()) {
+            thread_list.erase(it);
+            return;
+        }
+        ++it;
+    }
+    UNREACHABLE();
+}
 
-    return threads;
+std::vector<std::shared_ptr<Thread>> AddressArbiter::GetThreadsWaitingOnAddress(VAddr address) {
+    std::vector<std::shared_ptr<Thread>> result;
+    std::list<std::shared_ptr<Thread>>& thread_list = arb_threads[address];
+    auto it = thread_list.begin();
+    while (it != thread_list.end()) {
+        std::shared_ptr<Thread> current_thread = *it;
+        result.push_back(std::move(current_thread));
+        ++it;
+    }
+    return result;
 }
 } // namespace Kernel

src/core/hle/kernel/address_arbiter.h

@@ -4,10 +4,12 @@
 
 #pragma once
 
+#include <list>
+#include <memory>
+#include <unordered_map>
 #include <vector>
 
 #include "common/common_types.h"
-#include "core/hle/kernel/object.h"
 
 union ResultCode;
 
@@ -48,6 +50,9 @@ public:
     /// Waits on an address with a particular arbitration type.
     ResultCode WaitForAddress(VAddr address, ArbitrationType type, s32 value, s64 timeout_ns);
 
+    /// Removes a thread from the container and resets its address arbiter adress to 0
+    void HandleWakeupThread(std::shared_ptr<Thread> thread);
+
 private:
     /// Signals an address being waited on.
     ResultCode SignalToAddressOnly(VAddr address, s32 num_to_wake);
@@ -71,8 +76,20 @@ private:
     // Waits on the given address with a timeout in nanoseconds
     ResultCode WaitForAddressImpl(VAddr address, s64 timeout);
 
+    /// Wake up num_to_wake (or all) threads in a vector.
+    void WakeThreads(const std::vector<std::shared_ptr<Thread>>& waiting_threads, s32 num_to_wake);
+
+    /// Insert a thread into the address arbiter container
+    void InsertThread(std::shared_ptr<Thread> thread);
+
+    /// Removes a thread from the address arbiter container
+    void RemoveThread(std::shared_ptr<Thread> thread);
+
     // Gets the threads waiting on an address.
-    std::vector<std::shared_ptr<Thread>> GetThreadsWaitingOnAddress(VAddr address) const;
+    std::vector<std::shared_ptr<Thread>> GetThreadsWaitingOnAddress(VAddr address);
+
+    /// List of threads waiting for a address arbiter
+    std::unordered_map<VAddr, std::list<std::shared_ptr<Thread>>> arb_threads;
 
     Core::System& system;
 };

src/core/hle/kernel/client_port.cpp

@@ -8,7 +8,7 @@
 #include "core/hle/kernel/hle_ipc.h"
 #include "core/hle/kernel/object.h"
 #include "core/hle/kernel/server_port.h"
-#include "core/hle/kernel/server_session.h"
+#include "core/hle/kernel/session.h"
 
 namespace Kernel {
 
@@ -20,28 +20,23 @@ std::shared_ptr<ServerPort> ClientPort::GetServerPort() const {
 }
 
 ResultVal<std::shared_ptr<ClientSession>> ClientPort::Connect() {
-    // Note: Threads do not wait for the server endpoint to call
-    // AcceptSession before returning from this call.
-
     if (active_sessions >= max_sessions) {
         return ERR_MAX_CONNECTIONS_REACHED;
     }
     active_sessions++;
 
-    // Create a new session pair, let the created sessions inherit the parent port's HLE handler.
-    auto [server, client] =
-        ServerSession::CreateSessionPair(kernel, server_port->GetName(), SharedFrom(this));
+    auto [client, server] = Kernel::Session::Create(kernel, name);
 
     if (server_port->HasHLEHandler()) {
-        server_port->GetHLEHandler()->ClientConnected(server);
+        server_port->GetHLEHandler()->ClientConnected(std::move(server));
     } else {
-        server_port->AppendPendingSession(server);
+        server_port->AppendPendingSession(std::move(server));
     }
 
     // Wake the threads waiting on the ServerPort
     server_port->WakeupAllWaitingThreads();
 
-    return MakeResult(client);
+    return MakeResult(std::move(client));
 }
 
 void ClientPort::ConnectionClosed() {

src/core/hle/kernel/client_port.h

@@ -4,7 +4,9 @@
 
 #pragma once
 
+#include <memory>
 #include <string>
+
 #include "common/common_types.h"
 #include "core/hle/kernel/object.h"
 #include "core/hle/result.h"

src/core/hle/kernel/client_session.cpp

@@ -1,4 +1,4 @@
-// Copyright 2016 Citra Emulator Project
+// Copyright 2019 yuzu emulator team
 // Licensed under GPLv2 or any later version
 // Refer to the license.txt file included.
 
@@ -12,22 +12,44 @@
 
 namespace Kernel {
 
-ClientSession::ClientSession(KernelCore& kernel) : Object{kernel} {}
+ClientSession::ClientSession(KernelCore& kernel) : WaitObject{kernel} {}
+
 ClientSession::~ClientSession() {
     // This destructor will be called automatically when the last ClientSession handle is closed by
     // the emulated application.
-    if (auto server = parent->server.lock()) {
-        server->ClientDisconnected();
+    if (parent->Server()) {
+        parent->Server()->ClientDisconnected();
     }
 }
 
-ResultCode ClientSession::SendSyncRequest(Thread* thread, Memory::Memory& memory) {
-    // Signal the server session that new data is available
-    if (auto server = parent->server.lock()) {
-        return server->HandleSyncRequest(SharedFrom(thread), memory);
+bool ClientSession::ShouldWait(const Thread* thread) const {
+    UNIMPLEMENTED();
+    return {};
+}
+
+void ClientSession::Acquire(Thread* thread) {
+    UNIMPLEMENTED();
+}
+
+ResultVal<std::shared_ptr<ClientSession>> ClientSession::Create(KernelCore& kernel,
+                                                                std::shared_ptr<Session> parent,
+                                                                std::string name) {
+    std::shared_ptr<ClientSession> client_session{std::make_shared<ClientSession>(kernel)};
+
+    client_session->name = std::move(name);
+    client_session->parent = std::move(parent);
+
+    return MakeResult(std::move(client_session));
+}
+
+ResultCode ClientSession::SendSyncRequest(std::shared_ptr<Thread> thread, Memory::Memory& memory) {
+    // Keep ServerSession alive until we're done working with it.
+    if (!parent->Server()) {
+        return ERR_SESSION_CLOSED_BY_REMOTE;
     }
 
-    return ERR_SESSION_CLOSED_BY_REMOTE;
+    // Signal the server session that new data is available
+    return parent->Server()->HandleSyncRequest(std::move(thread), memory);
 }
 
 } // namespace Kernel

src/core/hle/kernel/client_session.h

@@ -1,4 +1,4 @@
-// Copyright 2016 Citra Emulator Project
+// Copyright 2019 yuzu emulator team
 // Licensed under GPLv2 or any later version
 // Refer to the license.txt file included.
 
@@ -6,7 +6,9 @@
 
 #include <memory>
 #include <string>
-#include "core/hle/kernel/object.h"
+
+#include "core/hle/kernel/wait_object.h"
+#include "core/hle/result.h"
 
 union ResultCode;
 
@@ -18,15 +20,14 @@ namespace Kernel {
 
 class KernelCore;
 class Session;
-class ServerSession;
 class Thread;
 
-class ClientSession final : public Object {
+class ClientSession final : public WaitObject {
 public:
     explicit ClientSession(KernelCore& kernel);
     ~ClientSession() override;
 
-    friend class ServerSession;
+    friend class Session;
 
     std::string GetTypeName() const override {
         return "ClientSession";
@@ -41,9 +42,17 @@ public:
         return HANDLE_TYPE;
     }
 
-    ResultCode SendSyncRequest(Thread* thread, Memory::Memory& memory);
+    ResultCode SendSyncRequest(std::shared_ptr<Thread> thread, Memory::Memory& memory);
+
+    bool ShouldWait(const Thread* thread) const override;
+
+    void Acquire(Thread* thread) override;
 
 private:
+    static ResultVal<std::shared_ptr<ClientSession>> Create(KernelCore& kernel,
+                                                            std::shared_ptr<Session> parent,
+                                                            std::string name = "Unknown");
+
     /// The parent session, which links to the server endpoint.
     std::shared_ptr<Session> parent;
 

src/core/hle/kernel/handle_table.h

@@ -6,6 +6,8 @@
 
 #include <array>
 #include <cstddef>
+#include <memory>
+
 #include "common/common_types.h"
 #include "core/hle/kernel/object.h"
 #include "core/hle/result.h"

src/core/hle/kernel/hle_ipc.cpp

@@ -74,6 +74,8 @@ std::shared_ptr<WritableEvent> HLERequestContext::SleepClientThread(
         thread->WakeAfterDelay(timeout);
     }
 
+    is_thread_waiting = true;
+
     return writable_event;
 }
 

src/core/hle/kernel/hle_ipc.h

@@ -264,6 +264,18 @@ public:
 
     std::string Description() const;
 
+    Thread& GetThread() {
+        return *thread;
+    }
+
+    const Thread& GetThread() const {
+        return *thread;
+    }
+
+    bool IsThreadWaiting() const {
+        return is_thread_waiting;
+    }
+
 private:
     void ParseCommandBuffer(const HandleTable& handle_table, u32_le* src_cmdbuf, bool incoming);
 
@@ -290,6 +302,7 @@ private:
     u32_le command{};
 
     std::vector<std::shared_ptr<SessionRequestHandler>> domain_request_handlers;
+    bool is_thread_waiting{};
 };
 
 } // namespace Kernel

src/core/hle/kernel/kernel.cpp

@@ -13,7 +13,6 @@
 #include "core/core.h"
 #include "core/core_timing.h"
 #include "core/core_timing_util.h"
-#include "core/hle/kernel/address_arbiter.h"
 #include "core/hle/kernel/client_port.h"
 #include "core/hle/kernel/errors.h"
 #include "core/hle/kernel/handle_table.h"
@@ -79,9 +78,9 @@ static void ThreadWakeupCallback(u64 thread_handle, [[maybe_unused]] s64 cycles_
         }
     }
 
-    if (thread->GetArbiterWaitAddress() != 0) {
-        ASSERT(thread->GetStatus() == ThreadStatus::WaitArb);
-        thread->SetArbiterWaitAddress(0);
+    if (thread->GetStatus() == ThreadStatus::WaitArb) {
+        auto& address_arbiter = thread->GetOwnerProcess()->GetAddressArbiter();
+        address_arbiter.HandleWakeupThread(thread);
     }
 
     if (resume) {

src/core/hle/kernel/kernel.h

@@ -4,6 +4,7 @@
 
 #pragma once
 
+#include <memory>
 #include <string>
 #include <unordered_map>
 #include <vector>

src/core/hle/kernel/mutex.cpp

@@ -2,6 +2,7 @@
 // Licensed under GPLv2 or any later version
 // Refer to the license.txt file included.
 
+#include <memory>
 #include <utility>
 #include <vector>
 

src/core/hle/kernel/object.cpp

@@ -27,6 +27,7 @@ bool Object::IsWaitable() const {
     case HandleType::ResourceLimit:
     case HandleType::ClientPort:
     case HandleType::ClientSession:
+    case HandleType::Session:
         return false;
     }
 

src/core/hle/kernel/object.h

@@ -29,6 +29,7 @@ enum class HandleType : u32 {
     ServerPort,
     ClientSession,
     ServerSession,
+    Session,
 };
 
 class Object : NonCopyable, public std::enable_shared_from_this<Object> {

src/core/hle/kernel/resource_limit.h

@@ -5,6 +5,8 @@
 #pragma once
 
 #include <array>
+#include <memory>
+
 #include "common/common_types.h"
 #include "core/hle/kernel/object.h"
 

src/core/hle/kernel/scheduler.cpp

@@ -458,7 +458,6 @@ void Scheduler::SwitchContext() {
         cpu_core.LoadContext(new_thread->GetContext());
         cpu_core.SetTlsAddress(new_thread->GetTLSAddress());
         cpu_core.SetTPIDR_EL0(new_thread->GetTPIDR_EL0());
-        cpu_core.ClearExclusiveState();
     } else {
         current_thread = nullptr;
         // Note: We do not reset the current process and current page table when idling because

src/core/hle/kernel/scheduler.h

@@ -4,11 +4,12 @@
 
 #pragma once
 
-#include <mutex>
+#include <atomic>
+#include <memory>
 #include <vector>
+
 #include "common/common_types.h"
 #include "common/multi_level_queue.h"
-#include "core/hle/kernel/object.h"
 #include "core/hle/kernel/thread.h"
 
 namespace Core {

src/core/hle/kernel/server_session.cpp

@@ -1,4 +1,4 @@
-// Copyright 2016 Citra Emulator Project
+// Copyright 2019 yuzu emulator team
 // Licensed under GPLv2 or any later version
 // Refer to the license.txt file included.
 
@@ -9,6 +9,7 @@
 #include "common/common_types.h"
 #include "common/logging/log.h"
 #include "core/core.h"
+#include "core/core_timing.h"
 #include "core/hle/ipc_helpers.h"
 #include "core/hle/kernel/client_port.h"
 #include "core/hle/kernel/client_session.h"
@@ -24,34 +25,29 @@
 namespace Kernel {
 
 ServerSession::ServerSession(KernelCore& kernel) : WaitObject{kernel} {}
-ServerSession::~ServerSession() {
-    // This destructor will be called automatically when the last ServerSession handle is closed by
-    // the emulated application.
-
-    // Decrease the port's connection count.
-    if (parent->port) {
-        parent->port->ConnectionClosed();
-    }
-}
+ServerSession::~ServerSession() = default;
 
 ResultVal<std::shared_ptr<ServerSession>> ServerSession::Create(KernelCore& kernel,
+                                                                std::shared_ptr<Session> parent,
                                                                 std::string name) {
-    std::shared_ptr<ServerSession> server_session = std::make_shared<ServerSession>(kernel);
+    std::shared_ptr<ServerSession> session{std::make_shared<ServerSession>(kernel)};
 
-    server_session->name = std::move(name);
-    server_session->parent = nullptr;
+    session->request_event = Core::Timing::CreateEvent(
+        name, [session](u64 userdata, s64 cycles_late) { session->CompleteSyncRequest(); });
+    session->name = std::move(name);
+    session->parent = std::move(parent);
 
-    return MakeResult(std::move(server_session));
+    return MakeResult(std::move(session));
 }
 
 bool ServerSession::ShouldWait(const Thread* thread) const {
-    // Wait if we have no pending requests, or if we're currently handling a request.
-    if (auto client = parent->client.lock()) {
-        return pending_requesting_threads.empty() || currently_handling != nullptr;
+    // Closed sessions should never wait, an error will be returned from svcReplyAndReceive.
+    if (!parent->Client()) {
+        return false;
     }
 
-    // Closed sessions should never wait, an error will be returned from svcReplyAndReceive.
-    return {};
+    // Wait if we have no pending requests, or if we're currently handling a request.
+    return pending_requesting_threads.empty() || currently_handling != nullptr;
 }
 
 void ServerSession::Acquire(Thread* thread) {
@@ -128,14 +124,21 @@ ResultCode ServerSession::HandleDomainSyncRequest(Kernel::HLERequestContext& con
     return RESULT_SUCCESS;
 }
 
-ResultCode ServerSession::HandleSyncRequest(std::shared_ptr<Thread> thread,
-                                            Memory::Memory& memory) {
-    // The ServerSession received a sync request, this means that there's new data available
-    // from its ClientSession, so wake up any threads that may be waiting on a svcReplyAndReceive or
-    // similar.
-    Kernel::HLERequestContext context(SharedFrom(this), thread);
-    u32* cmd_buf = (u32*)memory.GetPointer(thread->GetTLSAddress());
-    context.PopulateFromIncomingCommandBuffer(kernel.CurrentProcess()->GetHandleTable(), cmd_buf);
+ResultCode ServerSession::QueueSyncRequest(std::shared_ptr<Thread> thread, Memory::Memory& memory) {
+    u32* cmd_buf{reinterpret_cast<u32*>(memory.GetPointer(thread->GetTLSAddress()))};
+    std::shared_ptr<Kernel::HLERequestContext> context{
+        std::make_shared<Kernel::HLERequestContext>(SharedFrom(this), std::move(thread))};
+
+    context->PopulateFromIncomingCommandBuffer(kernel.CurrentProcess()->GetHandleTable(), cmd_buf);
+    request_queue.Push(std::move(context));
+
+    return RESULT_SUCCESS;
+}
+
+ResultCode ServerSession::CompleteSyncRequest() {
+    ASSERT(!request_queue.Empty());
+
+    auto& context = *request_queue.Front();
 
     ResultCode result = RESULT_SUCCESS;
     // If the session has been converted to a domain, handle the domain request
@@ -147,61 +150,27 @@ ResultCode ServerSession::HandleSyncRequest(std::shared_ptr<Thread> thread,
         result = hle_handler->HandleSyncRequest(context);
     }
 
-    if (thread->GetStatus() == ThreadStatus::Running) {
-        // Put the thread to sleep until the server replies, it will be awoken in
-        // svcReplyAndReceive for LLE servers.
-        thread->SetStatus(ThreadStatus::WaitIPC);
-
-        if (hle_handler != nullptr) {
-            // For HLE services, we put the request threads to sleep for a short duration to
-            // simulate IPC overhead, but only if the HLE handler didn't put the thread to sleep for
-            // other reasons like an async callback. The IPC overhead is needed to prevent
-            // starvation when a thread only does sync requests to HLE services while a
-            // lower-priority thread is waiting to run.
-
-            // This delay was approximated in a homebrew application by measuring the average time
-            // it takes for svcSendSyncRequest to return when performing the SetLcdForceBlack IPC
-            // request to the GSP:GPU service in a n3DS with firmware 11.6. The measured values have
-            // a high variance and vary between models.
-            static constexpr u64 IPCDelayNanoseconds = 39000;
-            thread->WakeAfterDelay(IPCDelayNanoseconds);
-        } else {
-            // Add the thread to the list of threads that have issued a sync request with this
-            // server.
-            pending_requesting_threads.push_back(std::move(thread));
-        }
-    }
-
-    // If this ServerSession does not have an HLE implementation, just wake up the threads waiting
-    // on it.
-    WakeupAllWaitingThreads();
-
-    // Handle scenario when ConvertToDomain command was issued, as we must do the conversion at the
-    // end of the command such that only commands following this one are handled as domains
     if (convert_to_domain) {
         ASSERT_MSG(IsSession(), "ServerSession is already a domain instance.");
         domain_request_handlers = {hle_handler};
         convert_to_domain = false;
     }
 
+    // Some service requests require the thread to block
+    if (!context.IsThreadWaiting()) {
+        context.GetThread().ResumeFromWait();
+        context.GetThread().SetWaitSynchronizationResult(result);
+    }
+
+    request_queue.Pop();
+
     return result;
 }
 
-ServerSession::SessionPair ServerSession::CreateSessionPair(KernelCore& kernel,
-                                                            const std::string& name,
-                                                            std::shared_ptr<ClientPort> port) {
-    auto server_session = ServerSession::Create(kernel, name + "_Server").Unwrap();
-    std::shared_ptr<ClientSession> client_session = std::make_shared<ClientSession>(kernel);
-    client_session->name = name + "_Client";
-
-    std::shared_ptr<Session> parent = std::make_shared<Session>();
-    parent->client = client_session;
-    parent->server = server_session;
-    parent->port = std::move(port);
-
-    client_session->parent = parent;
-    server_session->parent = parent;
-
-    return std::make_pair(std::move(server_session), std::move(client_session));
+ResultCode ServerSession::HandleSyncRequest(std::shared_ptr<Thread> thread,
+                                            Memory::Memory& memory) {
+    Core::System::GetInstance().CoreTiming().ScheduleEvent(20000, request_event, {});
+    return QueueSyncRequest(std::move(thread), memory);
 }
+
 } // namespace Kernel

src/core/hle/kernel/server_session.h

@@ -1,4 +1,4 @@
-// Copyright 2014 Citra Emulator Project
+// Copyright 2019 yuzu emulator team
 // Licensed under GPLv2 or any later version
 // Refer to the license.txt file included.
 
@@ -9,7 +9,7 @@
 #include <utility>
 #include <vector>
 
-#include "core/hle/kernel/object.h"
+#include "common/threadsafe_queue.h"
 #include "core/hle/kernel/wait_object.h"
 #include "core/hle/result.h"
 
@@ -17,13 +17,14 @@ namespace Memory {
 class Memory;
 }
 
+namespace Core::Timing {
+struct EventType;
+}
+
 namespace Kernel {
 
-class ClientPort;
-class ClientSession;
 class HLERequestContext;
 class KernelCore;
-class ServerSession;
 class Session;
 class SessionRequestHandler;
 class Thread;
@@ -45,6 +46,12 @@ public:
     explicit ServerSession(KernelCore& kernel);
     ~ServerSession() override;
 
+    friend class Session;
+
+    static ResultVal<std::shared_ptr<ServerSession>> Create(KernelCore& kernel,
+                                                            std::shared_ptr<Session> parent,
+                                                            std::string name = "Unknown");
+
     std::string GetTypeName() const override {
         return "ServerSession";
     }
@@ -66,18 +73,6 @@ public:
         return parent.get();
     }
 
-    using SessionPair = std::pair<std::shared_ptr<ServerSession>, std::shared_ptr<ClientSession>>;
-
-    /**
-     * Creates a pair of ServerSession and an associated ClientSession.
-     * @param kernel      The kernal instance to create the session pair under.
-     * @param name        Optional name of the ports.
-     * @param client_port Optional The ClientPort that spawned this session.
-     * @return The created session tuple
-     */
-    static SessionPair CreateSessionPair(KernelCore& kernel, const std::string& name = "Unknown",
-                                         std::shared_ptr<ClientPort> client_port = nullptr);
-
     /**
      * Sets the HLE handler for the session. This handler will be called to service IPC requests
      * instead of the regular IPC machinery. (The regular IPC machinery is currently not
@@ -128,15 +123,11 @@ public:
     }
 
 private:
-    /**
-     * Creates a server session. The server session can have an optional HLE handler,
-     * which will be invoked to handle the IPC requests that this session receives.
-     * @param kernel The kernel instance to create this server session under.
-     * @param name Optional name of the server session.
-     * @return The created server session
-     */
-    static ResultVal<std::shared_ptr<ServerSession>> Create(KernelCore& kernel,
-                                                            std::string name = "Unknown");
+    /// Queues a sync request from the emulated application.
+    ResultCode QueueSyncRequest(std::shared_ptr<Thread> thread, Memory::Memory& memory);
+
+    /// Completes a sync request from the emulated application.
+    ResultCode CompleteSyncRequest();
 
     /// Handles a SyncRequest to a domain, forwarding the request to the proper object or closing an
     /// object handle.
@@ -166,6 +157,12 @@ private:
 
     /// The name of this session (optional)
     std::string name;
+
+    /// Core timing event used to schedule the service request at some point in the future
+    std::shared_ptr<Core::Timing::EventType> request_event;
+
+    /// Queue of scheduled service requests
+    Common::MPSCQueue<std::shared_ptr<Kernel::HLERequestContext>> request_queue;
 };
 
 } // namespace Kernel

src/core/hle/kernel/session.cpp

@@ -1,12 +1,36 @@
-// Copyright 2015 Citra Emulator Project
+// Copyright 2019 yuzu emulator team
 // Licensed under GPLv2 or any later version
 // Refer to the license.txt file included.
 
+#include "common/assert.h"
+#include "core/hle/kernel/client_session.h"
+#include "core/hle/kernel/server_session.h"
 #include "core/hle/kernel/session.h"
-#include "core/hle/kernel/thread.h"
 
 namespace Kernel {
 
-Session::Session() {}
-Session::~Session() {}
+Session::Session(KernelCore& kernel) : WaitObject{kernel} {}
+Session::~Session() = default;
+
+Session::SessionPair Session::Create(KernelCore& kernel, std::string name) {
+    auto session{std::make_shared<Session>(kernel)};
+    auto client_session{Kernel::ClientSession::Create(kernel, session, name + "_Client").Unwrap()};
+    auto server_session{Kernel::ServerSession::Create(kernel, session, name + "_Server").Unwrap()};
+
+    session->name = std::move(name);
+    session->client = client_session;
+    session->server = server_session;
+
+    return std::make_pair(std::move(client_session), std::move(server_session));
+}
+
+bool Session::ShouldWait(const Thread* thread) const {
+    UNIMPLEMENTED();
+    return {};
+}
+
+void Session::Acquire(Thread* thread) {
+    UNIMPLEMENTED();
+}
+
 } // namespace Kernel

src/core/hle/kernel/session.h

@@ -1,27 +1,64 @@
-// Copyright 2018 yuzu emulator team
+// Copyright 2019 yuzu emulator team
 // Licensed under GPLv2 or any later version
 // Refer to the license.txt file included.
 
 #pragma once
 
-#include "core/hle/kernel/object.h"
+#include <memory>
+#include <string>
+#include <utility>
+
+#include "core/hle/kernel/wait_object.h"
 
 namespace Kernel {
 
 class ClientSession;
-class ClientPort;
 class ServerSession;
 
 /**
  * Parent structure to link the client and server endpoints of a session with their associated
- * client port. The client port need not exist, as is the case for portless sessions like the
- * FS File and Directory sessions. When one of the endpoints of a session is destroyed, its
- * corresponding field in this structure will be set to nullptr.
+ * client port.
  */
-class Session final {
+class Session final : public WaitObject {
 public:
-    std::weak_ptr<ClientSession> client; ///< The client endpoint of the session.
-    std::weak_ptr<ServerSession> server; ///< The server endpoint of the session.
-    std::shared_ptr<ClientPort> port; ///< The port that this session is associated with (optional).
+    explicit Session(KernelCore& kernel);
+    ~Session() override;
+
+    using SessionPair = std::pair<std::shared_ptr<ClientSession>, std::shared_ptr<ServerSession>>;
+
+    static SessionPair Create(KernelCore& kernel, std::string name = "Unknown");
+
+    std::string GetName() const override {
+        return name;
+    }
+
+    static constexpr HandleType HANDLE_TYPE = HandleType::Session;
+    HandleType GetHandleType() const override {
+        return HANDLE_TYPE;
+    }
+
+    bool ShouldWait(const Thread* thread) const override;
+
+    void Acquire(Thread* thread) override;
+
+    std::shared_ptr<ClientSession> Client() {
+        if (auto result{client.lock()}) {
+            return result;
+        }
+        return {};
+    }
+
+    std::shared_ptr<ServerSession> Server() {
+        if (auto result{server.lock()}) {
+            return result;
+        }
+        return {};
+    }
+
+private:
+    std::string name;
+    std::weak_ptr<ClientSession> client;
+    std::weak_ptr<ServerSession> server;
 };
+
 } // namespace Kernel

src/core/hle/kernel/shared_memory.h

@@ -6,7 +6,6 @@
 
 #include <memory>
 #include <string>
-#include <vector>
 
 #include "common/common_types.h"
 #include "core/hle/kernel/object.h"

src/core/hle/kernel/svc.cpp

@@ -381,11 +381,12 @@ static ResultCode SendSyncRequest(Core::System& system, Handle handle) {
 
     LOG_TRACE(Kernel_SVC, "called handle=0x{:08X}({})", handle, session->GetName());
 
-    system.PrepareReschedule();
+    auto thread = system.CurrentScheduler().GetCurrentThread();
+    thread->InvalidateWakeupCallback();
+    thread->SetStatus(ThreadStatus::WaitIPC);
+    system.PrepareReschedule(thread->GetProcessorID());
 
-    // TODO(Subv): svcSendSyncRequest should put the caller thread to sleep while the server
-    // responds and cause a reschedule.
-    return session->SendSyncRequest(system.CurrentScheduler().GetCurrentThread(), system.Memory());
+    return session->SendSyncRequest(SharedFrom(thread), system.Memory());
 }
 
 /// Get the ID for the specified thread.
@@ -1649,8 +1650,7 @@ static ResultCode WaitProcessWideKeyAtomic(Core::System& system, VAddr mutex_add
 }
 
 /// Signal process wide key
-static ResultCode SignalProcessWideKey(Core::System& system, VAddr condition_variable_addr,
-                                       s32 target) {
+static void SignalProcessWideKey(Core::System& system, VAddr condition_variable_addr, s32 target) {
     LOG_TRACE(Kernel_SVC, "called, condition_variable_addr=0x{:X}, target=0x{:08X}",
               condition_variable_addr, target);
 
@@ -1725,8 +1725,6 @@ static ResultCode SignalProcessWideKey(Core::System& system, VAddr condition_var
             system.PrepareReschedule(thread->GetProcessorID());
         }
     }
-
-    return RESULT_SUCCESS;
 }
 
 // Wait for an address (via Address Arbiter)
@@ -1780,6 +1778,17 @@ static ResultCode SignalToAddress(Core::System& system, VAddr address, u32 type,
     return address_arbiter.SignalToAddress(address, signal_type, value, num_to_wake);
 }
 
+static void KernelDebug([[maybe_unused]] Core::System& system,
+                        [[maybe_unused]] u32 kernel_debug_type, [[maybe_unused]] u64 param1,
+                        [[maybe_unused]] u64 param2, [[maybe_unused]] u64 param3) {
+    // Intentionally do nothing, as this does nothing in released kernel binaries.
+}
+
+static void ChangeKernelTraceState([[maybe_unused]] Core::System& system,
+                                   [[maybe_unused]] u32 trace_state) {
+    // Intentionally do nothing, as this does nothing in released kernel binaries.
+}
+
 /// This returns the total CPU ticks elapsed since the CPU was powered-on
 static u64 GetSystemTick(Core::System& system) {
     LOG_TRACE(Kernel_SVC, "called");
@@ -2417,8 +2426,8 @@ static const FunctionDef SVC_Table[] = {
     {0x39, nullptr, "Unknown"},
     {0x3A, nullptr, "Unknown"},
     {0x3B, nullptr, "Unknown"},
-    {0x3C, nullptr, "DumpInfo"},
-    {0x3D, nullptr, "DumpInfoNew"},
+    {0x3C, SvcWrap<KernelDebug>, "KernelDebug"},
+    {0x3D, SvcWrap<ChangeKernelTraceState>, "ChangeKernelTraceState"},
     {0x3E, nullptr, "Unknown"},
     {0x3F, nullptr, "Unknown"},
     {0x40, nullptr, "CreateSession"},

src/core/hle/kernel/svc_wrap.h

@@ -112,11 +112,6 @@ void SvcWrap(Core::System& system) {
     FuncReturn(system, retval);
 }
 
-template <ResultCode func(Core::System&, u64, s32)>
-void SvcWrap(Core::System& system) {
-    FuncReturn(system, func(system, Param(system, 0), static_cast<s32>(Param(system, 1))).raw);
-}
-
 template <ResultCode func(Core::System&, u64, u32)>
 void SvcWrap(Core::System& system) {
     FuncReturn(system, func(system, Param(system, 0), static_cast<u32>(Param(system, 1))).raw);
@@ -311,11 +306,27 @@ void SvcWrap(Core::System& system) {
     func(system);
 }
 
+template <void func(Core::System&, u32)>
+void SvcWrap(Core::System& system) {
+    func(system, static_cast<u32>(Param(system, 0)));
+}
+
+template <void func(Core::System&, u32, u64, u64, u64)>
+void SvcWrap(Core::System& system) {
+    func(system, static_cast<u32>(Param(system, 0)), Param(system, 1), Param(system, 2),
+         Param(system, 3));
+}
+
 template <void func(Core::System&, s64)>
 void SvcWrap(Core::System& system) {
     func(system, static_cast<s64>(Param(system, 0)));
 }
 
+template <void func(Core::System&, u64, s32)>
+void SvcWrap(Core::System& system) {
+    func(system, Param(system, 0), static_cast<s32>(Param(system, 1)));
+}
+
 template <void func(Core::System&, u64, u64)>
 void SvcWrap(Core::System& system) {
     func(system, Param(system, 0), Param(system, 1));

src/core/hle/kernel/transfer_memory.h

@@ -5,7 +5,6 @@
 #pragma once
 
 #include <memory>
-#include <vector>
 
 #include "core/hle/kernel/object.h"
 #include "core/hle/kernel/physical_memory.h"

src/core/hle/kernel/wait_object.h

@@ -4,8 +4,9 @@
 
 #pragma once
 
+#include <memory>
 #include <vector>
-#include <boost/smart_ptr/intrusive_ptr.hpp>
+
 #include "core/hle/kernel/object.h"
 
 namespace Kernel {

src/core/hle/kernel/writable_event.h

@@ -4,6 +4,8 @@
 
 #pragma once
 
+#include <memory>
+
 #include "core/hle/kernel/object.h"
 
 namespace Kernel {

src/core/hle/service/nfp/nfp.cpp

@@ -189,7 +189,7 @@ private:
         LOG_DEBUG(Service_NFP, "called");
 
         auto nfc_event = nfp_interface.GetNFCEvent();
-        if (!nfc_event->ShouldWait(Kernel::GetCurrentThread()) && !has_attached_handle) {
+        if (!nfc_event->ShouldWait(&ctx.GetThread()) && !has_attached_handle) {
             device_state = DeviceState::TagFound;
             nfc_event->Clear();
         }

src/core/hle/service/service.cpp

@@ -186,7 +186,7 @@ ResultCode ServiceFrameworkBase::HandleSyncRequest(Kernel::HLERequestContext& co
         UNIMPLEMENTED_MSG("command_type={}", static_cast<int>(context.GetCommandType()));
     }
 
-    context.WriteToOutgoingCommandBuffer(*Kernel::GetCurrentThread());
+    context.WriteToOutgoingCommandBuffer(context.GetThread());
 
     return RESULT_SUCCESS;
 }
@@ -201,7 +201,7 @@ void Init(std::shared_ptr<SM::ServiceManager>& sm, Core::System& system) {
     auto nv_flinger = std::make_shared<NVFlinger::NVFlinger>(system);
     system.GetFileSystemController().CreateFactories(*system.GetFilesystem(), false);
 
-    SM::ServiceManager::InstallInterfaces(sm);
+    SM::ServiceManager::InstallInterfaces(sm, system.Kernel());
 
     Account::InstallInterfaces(system);
     AM::InstallInterfaces(*sm, nv_flinger, system);

src/core/hle/service/sm/controller.cpp

@@ -30,10 +30,7 @@ void Controller::DuplicateSession(Kernel::HLERequestContext& ctx) {
 
     IPC::ResponseBuilder rb{ctx, 2, 0, 1, IPC::ResponseBuilder::Flags::AlwaysMoveHandles};
     rb.Push(RESULT_SUCCESS);
-    std::shared_ptr<Kernel::ClientSession> session{ctx.Session()->GetParent()->client};
-    rb.PushMoveObjects(session);
-
-    LOG_DEBUG(Service, "session={}", session->GetObjectId());
+    rb.PushMoveObjects(ctx.Session()->GetParent()->Client());
 }
 
 void Controller::DuplicateSessionEx(Kernel::HLERequestContext& ctx) {

src/core/hle/service/sm/sm.cpp

@@ -36,10 +36,11 @@ static ResultCode ValidateServiceName(const std::string& name) {
     return RESULT_SUCCESS;
 }
 
-void ServiceManager::InstallInterfaces(std::shared_ptr<ServiceManager> self) {
+void ServiceManager::InstallInterfaces(std::shared_ptr<ServiceManager> self,
+                                       Kernel::KernelCore& kernel) {
     ASSERT(self->sm_interface.expired());
 
-    auto sm = std::make_shared<SM>(self);
+    auto sm = std::make_shared<SM>(self, kernel);
     sm->InstallAsNamedPort();
     self->sm_interface = sm;
     self->controller_interface = std::make_unique<Controller>();
@@ -114,8 +115,6 @@ void SM::GetService(Kernel::HLERequestContext& ctx) {
 
     std::string name(name_buf.begin(), end);
 
-    // TODO(yuriks): Permission checks go here
-
     auto client_port = service_manager->GetServicePort(name);
     if (client_port.Failed()) {
         IPC::ResponseBuilder rb{ctx, 2};
@@ -127,14 +126,22 @@ void SM::GetService(Kernel::HLERequestContext& ctx) {
         return;
     }
 
-    auto session = client_port.Unwrap()->Connect();
-    ASSERT(session.Succeeded());
-    if (session.Succeeded()) {
-        LOG_DEBUG(Service_SM, "called service={} -> session={}", name, (*session)->GetObjectId());
-        IPC::ResponseBuilder rb{ctx, 2, 0, 1, IPC::ResponseBuilder::Flags::AlwaysMoveHandles};
-        rb.Push(session.Code());
-        rb.PushMoveObjects(std::move(session).Unwrap());
+    auto [client, server] = Kernel::Session::Create(kernel, name);
+
+    const auto& server_port = client_port.Unwrap()->GetServerPort();
+    if (server_port->GetHLEHandler()) {
+        server_port->GetHLEHandler()->ClientConnected(server);
+    } else {
+        server_port->AppendPendingSession(server);
     }
+
+    // Wake the threads waiting on the ServerPort
+    server_port->WakeupAllWaitingThreads();
+
+    LOG_DEBUG(Service_SM, "called service={} -> session={}", name, client->GetObjectId());
+    IPC::ResponseBuilder rb{ctx, 2, 0, 1, IPC::ResponseBuilder::Flags::AlwaysMoveHandles};
+    rb.Push(RESULT_SUCCESS);
+    rb.PushMoveObjects(std::move(client));
 }
 
 void SM::RegisterService(Kernel::HLERequestContext& ctx) {
@@ -178,8 +185,8 @@ void SM::UnregisterService(Kernel::HLERequestContext& ctx) {
     rb.Push(service_manager->UnregisterService(name));
 }
 
-SM::SM(std::shared_ptr<ServiceManager> service_manager)
-    : ServiceFramework("sm:", 4), service_manager(std::move(service_manager)) {
+SM::SM(std::shared_ptr<ServiceManager> service_manager, Kernel::KernelCore& kernel)
+    : ServiceFramework{"sm:", 4}, service_manager{std::move(service_manager)}, kernel{kernel} {
     static const FunctionInfo functions[] = {
         {0x00000000, &SM::Initialize, "Initialize"},
         {0x00000001, &SM::GetService, "GetService"},

src/core/hle/service/sm/sm.h

@@ -18,6 +18,7 @@
 namespace Kernel {
 class ClientPort;
 class ClientSession;
+class KernelCore;
 class ServerPort;
 class SessionRequestHandler;
 } // namespace Kernel
@@ -29,7 +30,7 @@ class Controller;
 /// Interface to "sm:" service
 class SM final : public ServiceFramework<SM> {
 public:
-    explicit SM(std::shared_ptr<ServiceManager> service_manager);
+    explicit SM(std::shared_ptr<ServiceManager> service_manager, Kernel::KernelCore& kernel);
     ~SM() override;
 
 private:
@@ -39,11 +40,12 @@ private:
     void UnregisterService(Kernel::HLERequestContext& ctx);
 
     std::shared_ptr<ServiceManager> service_manager;
+    Kernel::KernelCore& kernel;
 };
 
 class ServiceManager {
 public:
-    static void InstallInterfaces(std::shared_ptr<ServiceManager> self);
+    static void InstallInterfaces(std::shared_ptr<ServiceManager> self, Kernel::KernelCore& kernel);
 
     ServiceManager();
     ~ServiceManager();

src/core/telemetry_session.cpp

@@ -165,24 +165,20 @@ void TelemetrySession::AddInitialInfo(Loader::AppLoader& app_loader) {
     Telemetry::AppendOSInfo(field_collection);
 
     // Log user configuration information
-    AddField(Telemetry::FieldType::UserConfig, "Audio_SinkId", Settings::values.sink_id);
-    AddField(Telemetry::FieldType::UserConfig, "Audio_EnableAudioStretching",
-             Settings::values.enable_audio_stretching);
-    AddField(Telemetry::FieldType::UserConfig, "Core_UseMultiCore",
-             Settings::values.use_multi_core);
-    AddField(Telemetry::FieldType::UserConfig, "Renderer_ResolutionFactor",
-             Settings::values.resolution_factor);
-    AddField(Telemetry::FieldType::UserConfig, "Renderer_UseFrameLimit",
-             Settings::values.use_frame_limit);
-    AddField(Telemetry::FieldType::UserConfig, "Renderer_FrameLimit", Settings::values.frame_limit);
-    AddField(Telemetry::FieldType::UserConfig, "Renderer_UseDiskShaderCache",
-             Settings::values.use_disk_shader_cache);
-    AddField(Telemetry::FieldType::UserConfig, "Renderer_UseAccurateGpuEmulation",
+    constexpr auto field_type = Telemetry::FieldType::UserConfig;
+    AddField(field_type, "Audio_SinkId", Settings::values.sink_id);
+    AddField(field_type, "Audio_EnableAudioStretching", Settings::values.enable_audio_stretching);
+    AddField(field_type, "Core_UseMultiCore", Settings::values.use_multi_core);
+    AddField(field_type, "Renderer_Backend", "OpenGL");
+    AddField(field_type, "Renderer_ResolutionFactor", Settings::values.resolution_factor);
+    AddField(field_type, "Renderer_UseFrameLimit", Settings::values.use_frame_limit);
+    AddField(field_type, "Renderer_FrameLimit", Settings::values.frame_limit);
+    AddField(field_type, "Renderer_UseDiskShaderCache", Settings::values.use_disk_shader_cache);
+    AddField(field_type, "Renderer_UseAccurateGpuEmulation",
              Settings::values.use_accurate_gpu_emulation);
-    AddField(Telemetry::FieldType::UserConfig, "Renderer_UseAsynchronousGpuEmulation",
+    AddField(field_type, "Renderer_UseAsynchronousGpuEmulation",
              Settings::values.use_asynchronous_gpu_emulation);
-    AddField(Telemetry::FieldType::UserConfig, "System_UseDockedMode",
-             Settings::values.use_docked_mode);
+    AddField(field_type, "System_UseDockedMode", Settings::values.use_docked_mode);
 }
 
 bool TelemetrySession::SubmitTestcase() {

src/video_core/engines/maxwell_3d.h

@@ -310,6 +310,11 @@ public:
             }
         };
 
+        enum class DepthMode : u32 {
+            MinusOneToOne = 0,
+            ZeroToOne = 1,
+        };
+
         enum class PrimitiveTopology : u32 {
             Points = 0x0,
             Lines = 0x1,
@@ -491,6 +496,18 @@ public:
             INSERT_UNION_PADDING_WORDS(1);
         };
 
+        enum class TessellationPrimitive : u32 {
+            Isolines = 0,
+            Triangles = 1,
+            Quads = 2,
+        };
+
+        enum class TessellationSpacing : u32 {
+            Equal = 0,
+            FractionalOdd = 1,
+            FractionalEven = 2,
+        };
+
         struct RenderTargetConfig {
             u32 address_high;
             u32 address_low;
@@ -628,7 +645,19 @@ public:
                     };
                 } sync_info;
 
-                INSERT_UNION_PADDING_WORDS(0x11E);
+                INSERT_UNION_PADDING_WORDS(0x15);
+
+                union {
+                    BitField<0, 2, TessellationPrimitive> prim;
+                    BitField<4, 2, TessellationSpacing> spacing;
+                    BitField<8, 1, u32> cw;
+                    BitField<9, 1, u32> connected;
+                } tess_mode;
+
+                std::array<f32, 4> tess_level_outer;
+                std::array<f32, 2> tess_level_inner;
+
+                INSERT_UNION_PADDING_WORDS(0x102);
 
                 u32 tfb_enabled;
 
@@ -647,7 +676,7 @@ public:
                     u32 count;
                 } vertex_buffer;
 
-                INSERT_UNION_PADDING_WORDS(1);
+                DepthMode depth_mode;
 
                 float clear_color[4];
                 float clear_depth;
@@ -662,7 +691,9 @@ public:
                 u32 polygon_offset_line_enable;
                 u32 polygon_offset_fill_enable;
 
-                INSERT_UNION_PADDING_WORDS(0xD);
+                u32 patch_vertices;
+
+                INSERT_UNION_PADDING_WORDS(0xC);
 
                 std::array<ScissorTest, NumViewports> scissor_test;
 
@@ -1386,17 +1417,22 @@ ASSERT_REG_POSITION(upload, 0x60);
 ASSERT_REG_POSITION(exec_upload, 0x6C);
 ASSERT_REG_POSITION(data_upload, 0x6D);
 ASSERT_REG_POSITION(sync_info, 0xB2);
+ASSERT_REG_POSITION(tess_mode, 0xC8);
+ASSERT_REG_POSITION(tess_level_outer, 0xC9);
+ASSERT_REG_POSITION(tess_level_inner, 0xCD);
 ASSERT_REG_POSITION(tfb_enabled, 0x1D1);
 ASSERT_REG_POSITION(rt, 0x200);
 ASSERT_REG_POSITION(viewport_transform, 0x280);
 ASSERT_REG_POSITION(viewports, 0x300);
 ASSERT_REG_POSITION(vertex_buffer, 0x35D);
+ASSERT_REG_POSITION(depth_mode, 0x35F);
 ASSERT_REG_POSITION(clear_color[0], 0x360);
 ASSERT_REG_POSITION(clear_depth, 0x364);
 ASSERT_REG_POSITION(clear_stencil, 0x368);
 ASSERT_REG_POSITION(polygon_offset_point_enable, 0x370);
 ASSERT_REG_POSITION(polygon_offset_line_enable, 0x371);
 ASSERT_REG_POSITION(polygon_offset_fill_enable, 0x372);
+ASSERT_REG_POSITION(patch_vertices, 0x373);
 ASSERT_REG_POSITION(scissor_test, 0x380);
 ASSERT_REG_POSITION(stencil_back_func_ref, 0x3D5);
 ASSERT_REG_POSITION(stencil_back_func_mask, 0x3D6);

src/video_core/engines/shader_bytecode.h

@@ -98,10 +98,11 @@ union Attribute {
         BitField<20, 10, u64> immediate;
         BitField<22, 2, u64> element;
         BitField<24, 6, Index> index;
+        BitField<31, 1, u64> patch;
         BitField<47, 3, AttributeSize> size;
 
         bool IsPhysical() const {
-            return element == 0 && static_cast<u64>(index.Value()) == 0;
+            return patch == 0 && element == 0 && static_cast<u64>(index.Value()) == 0;
         }
     } fmt20;
 
@@ -383,6 +384,15 @@ enum class IsberdMode : u64 {
 
 enum class IsberdShift : u64 { None = 0, U16 = 1, B32 = 2 };
 
+enum class MembarType : u64 {
+    CTA = 0,
+    GL = 1,
+    SYS = 2,
+    VC = 3,
+};
+
+enum class MembarUnknown : u64 { Default = 0, IVALLD = 1, IVALLT = 2, IVALLTD = 3 };
+
 enum class HalfType : u64 {
     H0_H1 = 0,
     F32 = 1,
@@ -799,6 +809,12 @@ union Instruction {
         BitField<40, 1, u64> invert;
     } popc;
 
+    union {
+        BitField<41, 1, u64> sh;
+        BitField<40, 1, u64> invert;
+        BitField<48, 1, u64> is_signed;
+    } flo;
+
     union {
         BitField<39, 3, u64> pred;
         BitField<42, 1, u64> neg_pred;
@@ -1276,6 +1292,7 @@ union Instruction {
         BitField<50, 1, u64> dc_flag;
         BitField<51, 1, u64> aoffi_flag;
         BitField<52, 2, u64> component;
+        BitField<55, 1, u64> fp16_flag;
 
         bool UsesMiscMode(TextureMiscMode mode) const {
             switch (mode) {
@@ -1439,6 +1456,26 @@ union Instruction {
         }
     } tlds;
 
+    union {
+        BitField<28, 1, u64> is_array;
+        BitField<29, 2, TextureType> texture_type;
+        BitField<35, 1, u64> aoffi_flag;
+        BitField<49, 1, u64> nodep_flag;
+
+        bool UsesMiscMode(TextureMiscMode mode) const {
+            switch (mode) {
+            case TextureMiscMode::AOFFI:
+                return aoffi_flag != 0;
+            case TextureMiscMode::NODEP:
+                return nodep_flag != 0;
+            default:
+                break;
+            }
+            return false;
+        }
+
+    } txd;
+
     union {
         BitField<24, 2, StoreCacheManagement> cache_management;
         BitField<33, 3, ImageType> image_type;
@@ -1518,6 +1555,11 @@ union Instruction {
         BitField<47, 2, IsberdShift> shift;
     } isberd;
 
+    union {
+        BitField<8, 2, MembarType> type;
+        BitField<0, 2, MembarUnknown> unknown;
+    } membar;
+
     union {
         BitField<48, 1, u64> signed_a;
         BitField<38, 1, u64> is_byte_chunk_a;
@@ -1632,6 +1674,8 @@ public:
         TLD4S,  // Texture Load 4 with scalar / non - vec4 source / destinations
         TMML_B, // Texture Mip Map Level
         TMML,   // Texture Mip Map Level
+        TXD,    // Texture Gradient/Load with Derivates
+        TXD_B,  // Texture Gradient/Load with Derivates Bindless
         SUST,   // Surface Store
         SULD,   // Surface Load
         SUATOM, // Surface Atomic Operation
@@ -1640,6 +1684,7 @@ public:
         IPA,
         OUT_R, // Emit vertex/primitive
         ISBERD,
+        MEMBAR,
         VMAD,
         VSETP,
         FFMA_IMM, // Fused Multiply and Add
@@ -1664,6 +1709,9 @@ public:
         ISCADD_C, // Scale and Add
         ISCADD_R,
         ISCADD_IMM,
+        FLO_R,
+        FLO_C,
+        FLO_IMM,
         LEA_R1,
         LEA_R2,
         LEA_RZ,
@@ -1727,6 +1775,10 @@ public:
         SHR_C,
         SHR_R,
         SHR_IMM,
+        SHF_RIGHT_R,
+        SHF_RIGHT_IMM,
+        SHF_LEFT_R,
+        SHF_LEFT_IMM,
         FMNMX_C,
         FMNMX_R,
         FMNMX_IMM,
@@ -1894,7 +1946,7 @@ private:
             INST("111000100100----", Id::BRA, Type::Flow, "BRA"),
             INST("111000100101----", Id::BRX, Type::Flow, "BRX"),
             INST("1111000011111---", Id::SYNC, Type::Flow, "SYNC"),
-            INST("111000110100---", Id::BRK, Type::Flow, "BRK"),
+            INST("111000110100----", Id::BRK, Type::Flow, "BRK"),
             INST("111000110000----", Id::EXIT, Type::Flow, "EXIT"),
             INST("1111000011110---", Id::DEPBAR, Type::Synch, "DEPBAR"),
             INST("0101000011011---", Id::VOTE, Type::Warp, "VOTE"),
@@ -1921,9 +1973,11 @@ private:
             INST("1101-01---------", Id::TLDS, Type::Texture, "TLDS"),
             INST("110010----111---", Id::TLD4, Type::Texture, "TLD4"),
             INST("1101111011111---", Id::TLD4_B, Type::Texture, "TLD4_B"),
-            INST("1101111100------", Id::TLD4S, Type::Texture, "TLD4S"),
+            INST("11011111-0------", Id::TLD4S, Type::Texture, "TLD4S"),
             INST("110111110110----", Id::TMML_B, Type::Texture, "TMML_B"),
             INST("1101111101011---", Id::TMML, Type::Texture, "TMML"),
+            INST("11011110011110--", Id::TXD_B, Type::Texture, "TXD_B"),
+            INST("11011110001110--", Id::TXD, Type::Texture, "TXD"),
             INST("11101011001-----", Id::SUST, Type::Image, "SUST"),
             INST("11101011000-----", Id::SULD, Type::Image, "SULD"),
             INST("1110101000------", Id::SUATOM, Type::Image, "SUATOM_D"),
@@ -1931,6 +1985,7 @@ private:
             INST("11100000--------", Id::IPA, Type::Trivial, "IPA"),
             INST("1111101111100---", Id::OUT_R, Type::Trivial, "OUT_R"),
             INST("1110111111010---", Id::ISBERD, Type::Trivial, "ISBERD"),
+            INST("1110111110011---", Id::MEMBAR, Type::Trivial, "MEMBAR"),
             INST("01011111--------", Id::VMAD, Type::Video, "VMAD"),
             INST("0101000011110---", Id::VSETP, Type::Video, "VSETP"),
             INST("0011001-1-------", Id::FFMA_IMM, Type::Ffma, "FFMA_IMM"),
@@ -1965,6 +2020,9 @@ private:
             INST("010110110100----", Id::ICMP_R, Type::ArithmeticInteger, "ICMP_R"),
             INST("010010110100----", Id::ICMP_CR, Type::ArithmeticInteger, "ICMP_CR"),
             INST("0011011-0100----", Id::ICMP_IMM, Type::ArithmeticInteger, "ICMP_IMM"),
+            INST("0101110000110---", Id::FLO_R, Type::ArithmeticInteger, "FLO_R"),
+            INST("0100110000110---", Id::FLO_C, Type::ArithmeticInteger, "FLO_C"),
+            INST("0011100-00110---", Id::FLO_IMM, Type::ArithmeticInteger, "FLO_IMM"),
             INST("0101101111011---", Id::LEA_R2, Type::ArithmeticInteger, "LEA_R2"),
             INST("0101101111010---", Id::LEA_R1, Type::ArithmeticInteger, "LEA_R1"),
             INST("001101101101----", Id::LEA_IMM, Type::ArithmeticInteger, "LEA_IMM"),
@@ -2022,6 +2080,10 @@ private:
             INST("0100110000101---", Id::SHR_C, Type::Shift, "SHR_C"),
             INST("0101110000101---", Id::SHR_R, Type::Shift, "SHR_R"),
             INST("0011100-00101---", Id::SHR_IMM, Type::Shift, "SHR_IMM"),
+            INST("0101110011111---", Id::SHF_RIGHT_R, Type::Shift, "SHF_RIGHT_R"),
+            INST("0011100-11111---", Id::SHF_RIGHT_IMM, Type::Shift, "SHF_RIGHT_IMM"),
+            INST("0101101111111---", Id::SHF_LEFT_R, Type::Shift, "SHF_LEFT_R"),
+            INST("0011011-11111---", Id::SHF_LEFT_IMM, Type::Shift, "SHF_LEFT_IMM"),
             INST("0100110011100---", Id::I2I_C, Type::Conversion, "I2I_C"),
             INST("0101110011100---", Id::I2I_R, Type::Conversion, "I2I_R"),
             INST("0011101-11100---", Id::I2I_IMM, Type::Conversion, "I2I_IMM"),

src/video_core/rasterizer_accelerated.cpp

@@ -5,6 +5,7 @@
 #include <mutex>
 
 #include <boost/icl/interval_map.hpp>
+#include <boost/range/iterator_range.hpp>
 
 #include "common/assert.h"
 #include "common/common_types.h"

src/video_core/renderer_opengl/gl_device.cpp

@@ -5,6 +5,7 @@
 #include <algorithm>
 #include <array>
 #include <cstddef>
+#include <cstring>
 #include <optional>
 #include <vector>
 
@@ -134,11 +135,13 @@ std::array<Device::BaseBindings, Tegra::Engines::MaxShaderTypes> BuildBaseBindin
 
 Device::Device() : base_bindings{BuildBaseBindings()} {
     const std::string_view vendor = reinterpret_cast<const char*>(glGetString(GL_VENDOR));
+    const auto renderer = reinterpret_cast<const char*>(glGetString(GL_RENDERER));
     const std::vector extensions = GetExtensions();
 
     const bool is_nvidia = vendor == "NVIDIA Corporation";
     const bool is_amd = vendor == "ATI Technologies Inc.";
     const bool is_intel = vendor == "Intel";
+    const bool is_intel_proprietary = is_intel && std::strstr(renderer, "Mesa") == nullptr;
 
     uniform_buffer_alignment = GetInteger<std::size_t>(GL_UNIFORM_BUFFER_OFFSET_ALIGNMENT);
     shader_storage_alignment = GetInteger<std::size_t>(GL_SHADER_STORAGE_BUFFER_OFFSET_ALIGNMENT);
@@ -152,7 +155,7 @@ Device::Device() : base_bindings{BuildBaseBindings()} {
     has_variable_aoffi = TestVariableAoffi();
     has_component_indexing_bug = is_amd;
     has_precise_bug = TestPreciseBug();
-    has_broken_compute = is_intel;
+    has_broken_compute = is_intel_proprietary;
     has_fast_buffer_sub_data = is_nvidia;
 
     LOG_INFO(Render_OpenGL, "Renderer_VariableAOFFI: {}", has_variable_aoffi);

src/video_core/renderer_opengl/gl_framebuffer_cache.cpp

@@ -3,9 +3,12 @@
 // Refer to the license.txt file included.
 
 #include <tuple>
+#include <unordered_map>
+#include <utility>
 
-#include "common/cityhash.h"
-#include "common/scope_exit.h"
+#include <glad/glad.h>
+
+#include "common/common_types.h"
 #include "video_core/engines/maxwell_3d.h"
 #include "video_core/renderer_opengl/gl_framebuffer_cache.h"
 #include "video_core/renderer_opengl/gl_state.h"
@@ -13,6 +16,7 @@
 namespace OpenGL {
 
 using Maxwell = Tegra::Engines::Maxwell3D::Regs;
+using VideoCore::Surface::SurfaceType;
 
 FramebufferCacheOpenGL::FramebufferCacheOpenGL() = default;
 
@@ -35,36 +39,49 @@ OGLFramebuffer FramebufferCacheOpenGL::CreateFramebuffer(const FramebufferCacheK
     local_state.draw.draw_framebuffer = framebuffer.handle;
     local_state.ApplyFramebufferState();
 
-    for (std::size_t index = 0; index < Maxwell::NumRenderTargets; ++index) {
-        if (key.colors[index]) {
-            key.colors[index]->Attach(GL_COLOR_ATTACHMENT0 + static_cast<GLenum>(index),
-                                      GL_DRAW_FRAMEBUFFER);
-        }
-    }
-    if (key.colors_count) {
-        glDrawBuffers(key.colors_count, key.color_attachments.data());
-    } else {
-        glDrawBuffer(GL_NONE);
+    if (key.zeta) {
+        const bool stencil = key.zeta->GetSurfaceParams().type == SurfaceType::DepthStencil;
+        const GLenum attach_target = stencil ? GL_DEPTH_STENCIL_ATTACHMENT : GL_DEPTH_ATTACHMENT;
+        key.zeta->Attach(attach_target, GL_DRAW_FRAMEBUFFER);
     }
 
-    if (key.zeta) {
-        key.zeta->Attach(key.stencil_enable ? GL_DEPTH_STENCIL_ATTACHMENT : GL_DEPTH_ATTACHMENT,
-                         GL_DRAW_FRAMEBUFFER);
+    std::size_t num_buffers = 0;
+    std::array<GLenum, Maxwell::NumRenderTargets> targets;
+
+    for (std::size_t index = 0; index < Maxwell::NumRenderTargets; ++index) {
+        if (!key.colors[index]) {
+            targets[index] = GL_NONE;
+            continue;
+        }
+        const GLenum attach_target = GL_COLOR_ATTACHMENT0 + static_cast<GLenum>(index);
+        key.colors[index]->Attach(attach_target, GL_DRAW_FRAMEBUFFER);
+
+        const u32 attachment = (key.color_attachments >> (BitsPerAttachment * index)) & 0b1111;
+        targets[index] = GL_COLOR_ATTACHMENT0 + attachment;
+        num_buffers = index + 1;
+    }
+
+    if (num_buffers > 0) {
+        glDrawBuffers(static_cast<GLsizei>(num_buffers), std::data(targets));
+    } else {
+        glDrawBuffer(GL_NONE);
     }
 
     return framebuffer;
 }
 
-std::size_t FramebufferCacheKey::Hash() const {
-    static_assert(sizeof(*this) % sizeof(u64) == 0, "Unaligned struct");
-    return static_cast<std::size_t>(
-        Common::CityHash64(reinterpret_cast<const char*>(this), sizeof(*this)));
+std::size_t FramebufferCacheKey::Hash() const noexcept {
+    std::size_t hash = std::hash<View>{}(zeta);
+    for (const auto& color : colors) {
+        hash ^= std::hash<View>{}(color);
+    }
+    hash ^= static_cast<std::size_t>(color_attachments) << 16;
+    return hash;
 }
 
-bool FramebufferCacheKey::operator==(const FramebufferCacheKey& rhs) const {
-    return std::tie(stencil_enable, colors_count, color_attachments, colors, zeta) ==
-           std::tie(rhs.stencil_enable, rhs.colors_count, rhs.color_attachments, rhs.colors,
-                    rhs.zeta);
+bool FramebufferCacheKey::operator==(const FramebufferCacheKey& rhs) const noexcept {
+    return std::tie(colors, zeta, color_attachments) ==
+           std::tie(rhs.colors, rhs.zeta, rhs.color_attachments);
 }
 
 } // namespace OpenGL

src/video_core/renderer_opengl/gl_framebuffer_cache.h

@@ -18,21 +18,24 @@
 
 namespace OpenGL {
 
-struct alignas(sizeof(u64)) FramebufferCacheKey {
-    bool stencil_enable = false;
-    u16 colors_count = 0;
+constexpr std::size_t BitsPerAttachment = 4;
 
-    std::array<GLenum, Tegra::Engines::Maxwell3D::Regs::NumRenderTargets> color_attachments{};
-    std::array<View, Tegra::Engines::Maxwell3D::Regs::NumRenderTargets> colors;
+struct FramebufferCacheKey {
     View zeta;
+    std::array<View, Tegra::Engines::Maxwell3D::Regs::NumRenderTargets> colors;
+    u32 color_attachments = 0;
 
-    std::size_t Hash() const;
+    std::size_t Hash() const noexcept;
 
-    bool operator==(const FramebufferCacheKey& rhs) const;
+    bool operator==(const FramebufferCacheKey& rhs) const noexcept;
 
-    bool operator!=(const FramebufferCacheKey& rhs) const {
+    bool operator!=(const FramebufferCacheKey& rhs) const noexcept {
         return !operator==(rhs);
     }
+
+    void SetAttachment(std::size_t index, u32 attachment) {
+        color_attachments |= attachment << (BitsPerAttachment * index);
+    }
 };
 
 } // namespace OpenGL

src/video_core/renderer_opengl/gl_rasterizer.cpp

@@ -93,7 +93,6 @@ RasterizerOpenGL::RasterizerOpenGL(Core::System& system, Core::Frontend::EmuWind
     shader_program_manager = std::make_unique<GLShader::ProgramManager>();
     state.draw.shader_program = 0;
     state.Apply();
-    clear_framebuffer.Create();
 
     LOG_DEBUG(Render_OpenGL, "Sync fixed function OpenGL state here");
     CheckExtensions();
@@ -278,6 +277,14 @@ void RasterizerOpenGL::SetupShaders(GLenum primitive_mode) {
             continue;
         }
 
+        // Currently this stages are not supported in the OpenGL backend.
+        // Todo(Blinkhawk): Port tesselation shaders from Vulkan to OpenGL
+        if (program == Maxwell::ShaderProgram::TesselationControl) {
+            continue;
+        } else if (program == Maxwell::ShaderProgram::TesselationEval) {
+            continue;
+        }
+
         Shader shader{shader_cache.GetStageProgram(program)};
 
         // Stage indices are 0 - 5
@@ -373,78 +380,58 @@ void RasterizerOpenGL::ConfigureFramebuffers() {
     UNIMPLEMENTED_IF(regs.rt_separate_frag_data == 0);
 
     // Bind the framebuffer surfaces
-    FramebufferCacheKey fbkey;
-    for (std::size_t index = 0; index < Maxwell::NumRenderTargets; ++index) {
+    FramebufferCacheKey key;
+    const auto colors_count = static_cast<std::size_t>(regs.rt_control.count);
+    for (std::size_t index = 0; index < colors_count; ++index) {
         View color_surface{texture_cache.GetColorBufferSurface(index, true)};
-
-        if (color_surface) {
-            // Assume that a surface will be written to if it is used as a framebuffer, even
-            // if the shader doesn't actually write to it.
-            texture_cache.MarkColorBufferInUse(index);
+        if (!color_surface) {
+            continue;
         }
+        // Assume that a surface will be written to if it is used as a framebuffer, even
+        // if the shader doesn't actually write to it.
+        texture_cache.MarkColorBufferInUse(index);
 
-        fbkey.color_attachments[index] = GL_COLOR_ATTACHMENT0 + regs.rt_control.GetMap(index);
-        fbkey.colors[index] = std::move(color_surface);
+        key.SetAttachment(index, regs.rt_control.GetMap(index));
+        key.colors[index] = std::move(color_surface);
     }
-    fbkey.colors_count = static_cast<u16>(regs.rt_control.count);
 
     if (depth_surface) {
         // Assume that a surface will be written to if it is used as a framebuffer, even if
         // the shader doesn't actually write to it.
         texture_cache.MarkDepthBufferInUse();
-
-        fbkey.stencil_enable = depth_surface->GetSurfaceParams().type == SurfaceType::DepthStencil;
-        fbkey.zeta = std::move(depth_surface);
+        key.zeta = std::move(depth_surface);
     }
 
     texture_cache.GuardRenderTargets(false);
 
-    state.draw.draw_framebuffer = framebuffer_cache.GetFramebuffer(fbkey);
+    state.draw.draw_framebuffer = framebuffer_cache.GetFramebuffer(key);
     SyncViewport(state);
 }
 
 void RasterizerOpenGL::ConfigureClearFramebuffer(OpenGLState& current_state, bool using_color_fb,
                                                  bool using_depth_fb, bool using_stencil_fb) {
+    using VideoCore::Surface::SurfaceType;
+
     auto& gpu = system.GPU().Maxwell3D();
     const auto& regs = gpu.regs;
 
     texture_cache.GuardRenderTargets(true);
-    View color_surface{};
+    View color_surface;
     if (using_color_fb) {
         color_surface = texture_cache.GetColorBufferSurface(regs.clear_buffers.RT, false);
     }
-    View depth_surface{};
+    View depth_surface;
     if (using_depth_fb || using_stencil_fb) {
         depth_surface = texture_cache.GetDepthBufferSurface(false);
     }
     texture_cache.GuardRenderTargets(false);
 
-    current_state.draw.draw_framebuffer = clear_framebuffer.handle;
+    FramebufferCacheKey key;
+    key.colors[0] = color_surface;
+    key.zeta = depth_surface;
+
+    current_state.draw.draw_framebuffer = framebuffer_cache.GetFramebuffer(key);
     current_state.ApplyFramebufferState();
-
-    if (color_surface) {
-        color_surface->Attach(GL_COLOR_ATTACHMENT0, GL_DRAW_FRAMEBUFFER);
-    } else {
-        glFramebufferTexture2D(GL_DRAW_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, 0, 0);
-    }
-
-    if (depth_surface) {
-        const auto& params = depth_surface->GetSurfaceParams();
-        switch (params.type) {
-        case VideoCore::Surface::SurfaceType::Depth:
-            depth_surface->Attach(GL_DEPTH_ATTACHMENT, GL_DRAW_FRAMEBUFFER);
-            glFramebufferTexture2D(GL_DRAW_FRAMEBUFFER, GL_STENCIL_ATTACHMENT, GL_TEXTURE_2D, 0, 0);
-            break;
-        case VideoCore::Surface::SurfaceType::DepthStencil:
-            depth_surface->Attach(GL_DEPTH_STENCIL_ATTACHMENT, GL_DRAW_FRAMEBUFFER);
-            break;
-        default:
-            UNIMPLEMENTED();
-        }
-    } else {
-        glFramebufferTexture2D(GL_DRAW_FRAMEBUFFER, GL_DEPTH_STENCIL_ATTACHMENT, GL_TEXTURE_2D, 0,
-                               0);
-    }
 }
 
 void RasterizerOpenGL::Clear() {
@@ -1049,6 +1036,10 @@ void RasterizerOpenGL::SyncViewport(OpenGLState& current_state) {
         flip_y = !flip_y;
     }
     state.clip_control.origin = flip_y ? GL_UPPER_LEFT : GL_LOWER_LEFT;
+    state.clip_control.depth_mode =
+        regs.depth_mode == Tegra::Engines::Maxwell3D::Regs::DepthMode::ZeroToOne
+            ? GL_ZERO_TO_ONE
+            : GL_NEGATIVE_ONE_TO_ONE;
 }
 
 void RasterizerOpenGL::SyncClipEnabled(

src/video_core/renderer_opengl/gl_rasterizer.h

@@ -223,8 +223,6 @@ private:
 
     enum class AccelDraw { Disabled, Arrays, Indexed };
     AccelDraw accelerate_draw = AccelDraw::Disabled;
-
-    OGLFramebuffer clear_framebuffer;
 };
 
 } // namespace OpenGL

src/video_core/renderer_opengl/gl_shader_cache.cpp

@@ -281,11 +281,11 @@ CachedProgram BuildShader(const Device& device, u64 unique_identifier, ShaderTyp
         if (variant.shared_memory_size > 0) {
             // TODO(Rodrigo): We should divide by four here, but having a larger shared memory pool
             // avoids out of bound stores. Find out why shared memory size is being invalid.
-            source += fmt::format("shared uint smem[{}];", variant.shared_memory_size);
+            source += fmt::format("shared uint smem[{}];\n", variant.shared_memory_size);
         }
 
         if (variant.local_memory_size > 0) {
-            source += fmt::format("#define LOCAL_MEMORY_SIZE {}",
+            source += fmt::format("#define LOCAL_MEMORY_SIZE {}\n",
                                   Common::AlignUp(variant.local_memory_size, 4) / 4);
         }
     }

src/video_core/renderer_opengl/gl_shader_decompiler.cpp

@@ -49,8 +49,9 @@ class ExprDecompiler;
 enum class Type { Void, Bool, Bool2, Float, Int, Uint, HalfFloat };
 
 struct TextureAoffi {};
+struct TextureDerivates {};
 using TextureArgument = std::pair<Type, Node>;
-using TextureIR = std::variant<TextureAoffi, TextureArgument>;
+using TextureIR = std::variant<TextureAoffi, TextureDerivates, TextureArgument>;
 
 constexpr u32 MAX_CONSTBUFFER_ELEMENTS =
     static_cast<u32>(Maxwell::MaxConstBufferSize) / (4 * sizeof(float));
@@ -398,6 +399,7 @@ public:
         DeclareConstantBuffers();
         DeclareGlobalMemory();
         DeclareSamplers();
+        DeclareImages();
         DeclarePhysicalAttributeReader();
 
         code.AddLine("void execute_{}() {{", suffix);
@@ -1075,7 +1077,7 @@ private:
     }
 
     std::string GenerateTexture(Operation operation, const std::string& function_suffix,
-                                const std::vector<TextureIR>& extras) {
+                                const std::vector<TextureIR>& extras, bool sepparate_dc = false) {
         constexpr std::array coord_constructors = {"float", "vec2", "vec3", "vec4"};
 
         const auto meta = std::get_if<MetaTexture>(&operation.GetMeta());
@@ -1090,7 +1092,8 @@ private:
             expr += "Offset";
         }
         expr += '(' + GetSampler(meta->sampler) + ", ";
-        expr += coord_constructors.at(count + (has_array ? 1 : 0) + (has_shadow ? 1 : 0) - 1);
+        expr += coord_constructors.at(count + (has_array ? 1 : 0) +
+                                      (has_shadow && !sepparate_dc ? 1 : 0) - 1);
         expr += '(';
         for (std::size_t i = 0; i < count; ++i) {
             expr += Visit(operation[i]).AsFloat();
@@ -1103,15 +1106,22 @@ private:
             expr += ", float(" + Visit(meta->array).AsInt() + ')';
         }
         if (has_shadow) {
-            expr += ", " + Visit(meta->depth_compare).AsFloat();
+            if (sepparate_dc) {
+                expr += "), " + Visit(meta->depth_compare).AsFloat();
+            } else {
+                expr += ", " + Visit(meta->depth_compare).AsFloat() + ')';
+            }
+        } else {
+            expr += ')';
         }
-        expr += ')';
 
         for (const auto& variant : extras) {
             if (const auto argument = std::get_if<TextureArgument>(&variant)) {
                 expr += GenerateTextureArgument(*argument);
             } else if (std::holds_alternative<TextureAoffi>(variant)) {
                 expr += GenerateTextureAoffi(meta->aoffi);
+            } else if (std::holds_alternative<TextureDerivates>(variant)) {
+                expr += GenerateTextureDerivates(meta->derivates);
             } else {
                 UNREACHABLE();
             }
@@ -1181,6 +1191,36 @@ private:
         return expr;
     }
 
+    std::string GenerateTextureDerivates(const std::vector<Node>& derivates) {
+        if (derivates.empty()) {
+            return {};
+        }
+        constexpr std::array coord_constructors = {"float", "vec2", "vec3"};
+        std::string expr = ", ";
+        const std::size_t components = derivates.size() / 2;
+        std::string dx = coord_constructors.at(components - 1);
+        std::string dy = coord_constructors.at(components - 1);
+        dx += '(';
+        dy += '(';
+
+        for (std::size_t index = 0; index < components; ++index) {
+            const auto operand_x{derivates.at(index * 2)};
+            const auto operand_y{derivates.at(index * 2 + 1)};
+            dx += Visit(operand_x).AsFloat();
+            dy += Visit(operand_y).AsFloat();
+
+            if (index + 1 < components) {
+                dx += ", ";
+                dy += ", ";
+            }
+        }
+        dx += ')';
+        dy += ')';
+        expr += dx + ", " + dy;
+
+        return expr;
+    }
+
     std::string BuildIntegerCoordinates(Operation operation) {
         constexpr std::array constructors{"int(", "ivec2(", "ivec3(", "ivec4("};
         const std::size_t coords_count{operation.GetOperandsCount()};
@@ -1450,6 +1490,11 @@ private:
         return GenerateUnary(operation, "bitCount", type, type);
     }
 
+    template <Type type>
+    Expression BitMSB(Operation operation) {
+        return GenerateUnary(operation, "findMSB", type, type);
+    }
+
     Expression HNegate(Operation operation) {
         const auto GetNegate = [&](std::size_t index) {
             return VisitOperand(operation, index).AsBool() + " ? -1 : 1";
@@ -1668,10 +1713,17 @@ private:
         ASSERT(meta);
 
         const auto type = meta->sampler.IsShadow() ? Type::Float : Type::Int;
-        return {GenerateTexture(operation, "Gather",
-                                {TextureAoffi{}, TextureArgument{type, meta->component}}) +
-                    GetSwizzle(meta->element),
-                Type::Float};
+        if (meta->sampler.IsShadow()) {
+            return {GenerateTexture(operation, "Gather", {TextureAoffi{}}, true) +
+                        GetSwizzle(meta->element),
+                    Type::Float};
+        } else {
+            return {GenerateTexture(operation, "Gather",
+                                    {TextureAoffi{}, TextureArgument{type, meta->component}},
+                                    false) +
+                        GetSwizzle(meta->element),
+                    Type::Float};
+        }
     }
 
     Expression TextureQueryDimensions(Operation operation) {
@@ -1738,6 +1790,14 @@ private:
         return {std::move(expr), Type::Float};
     }
 
+    Expression TextureGradient(Operation operation) {
+        const auto meta = std::get_if<MetaTexture>(&operation.GetMeta());
+        ASSERT(meta);
+
+        std::string expr = GenerateTexture(operation, "Grad", {TextureDerivates{}, TextureAoffi{}});
+        return {std::move(expr) + GetSwizzle(meta->element), Type::Float};
+    }
+
     Expression ImageLoad(Operation operation) {
         if (!device.HasImageLoadFormatted()) {
             LOG_ERROR(Render_OpenGL,
@@ -1869,6 +1929,10 @@ private:
         return {};
     }
 
+    Expression InvocationId(Operation operation) {
+        return {"gl_InvocationID", Type::Int};
+    }
+
     Expression YNegate(Operation operation) {
         return {"y_direction", Type::Float};
     }
@@ -1942,6 +2006,11 @@ private:
         return {fmt::format("readInvocationARB({}, {})", value, index), Type::Float};
     }
 
+    Expression MemoryBarrierGL(Operation) {
+        code.AddLine("memoryBarrier();");
+        return {};
+    }
+
     struct Func final {
         Func() = delete;
         ~Func() = delete;
@@ -2003,6 +2072,7 @@ private:
         &GLSLDecompiler::BitfieldInsert<Type::Int>,
         &GLSLDecompiler::BitfieldExtract<Type::Int>,
         &GLSLDecompiler::BitCount<Type::Int>,
+        &GLSLDecompiler::BitMSB<Type::Int>,
 
         &GLSLDecompiler::Add<Type::Uint>,
         &GLSLDecompiler::Mul<Type::Uint>,
@@ -2021,6 +2091,7 @@ private:
         &GLSLDecompiler::BitfieldInsert<Type::Uint>,
         &GLSLDecompiler::BitfieldExtract<Type::Uint>,
         &GLSLDecompiler::BitCount<Type::Uint>,
+        &GLSLDecompiler::BitMSB<Type::Uint>,
 
         &GLSLDecompiler::Add<Type::HalfFloat>,
         &GLSLDecompiler::Mul<Type::HalfFloat>,
@@ -2084,6 +2155,7 @@ private:
         &GLSLDecompiler::TextureQueryDimensions,
         &GLSLDecompiler::TextureQueryLod,
         &GLSLDecompiler::TexelFetch,
+        &GLSLDecompiler::TextureGradient,
 
         &GLSLDecompiler::ImageLoad,
         &GLSLDecompiler::ImageStore,
@@ -2104,6 +2176,7 @@ private:
         &GLSLDecompiler::EmitVertex,
         &GLSLDecompiler::EndPrimitive,
 
+        &GLSLDecompiler::InvocationId,
         &GLSLDecompiler::YNegate,
         &GLSLDecompiler::LocalInvocationId<0>,
         &GLSLDecompiler::LocalInvocationId<1>,
@@ -2119,6 +2192,8 @@ private:
 
         &GLSLDecompiler::ThreadId,
         &GLSLDecompiler::ShuffleIndexed,
+
+        &GLSLDecompiler::MemoryBarrierGL,
     };
     static_assert(operation_decompilers.size() == static_cast<std::size_t>(OperationCode::Amount));
 

src/video_core/renderer_opengl/gl_state.cpp

@@ -411,8 +411,9 @@ void OpenGLState::ApplyAlphaTest() {
 }
 
 void OpenGLState::ApplyClipControl() {
-    if (UpdateValue(cur_state.clip_control.origin, clip_control.origin)) {
-        glClipControl(clip_control.origin, GL_NEGATIVE_ONE_TO_ONE);
+    if (UpdateTie(std::tie(cur_state.clip_control.origin, cur_state.clip_control.depth_mode),
+                  std::tie(clip_control.origin, clip_control.depth_mode))) {
+        glClipControl(clip_control.origin, clip_control.depth_mode);
     }
 }
 

src/video_core/renderer_opengl/gl_state.h

@@ -150,6 +150,7 @@ public:
 
     struct {
         GLenum origin = GL_LOWER_LEFT;
+        GLenum depth_mode = GL_NEGATIVE_ONE_TO_ONE;
     } clip_control;
 
     OpenGLState();

src/video_core/renderer_opengl/renderer_opengl.cpp

@@ -24,19 +24,21 @@
 
 namespace OpenGL {
 
-static const char vertex_shader[] = R"(
-#version 150 core
+namespace {
 
-in vec2 vert_position;
-in vec2 vert_tex_coord;
-out vec2 frag_tex_coord;
+constexpr char vertex_shader[] = R"(
+#version 430 core
+
+layout (location = 0) in vec2 vert_position;
+layout (location = 1) in vec2 vert_tex_coord;
+layout (location = 0) out vec2 frag_tex_coord;
 
 // This is a truncated 3x3 matrix for 2D transformations:
 // The upper-left 2x2 submatrix performs scaling/rotation/mirroring.
 // The third column performs translation.
 // The third row could be used for projection, which we don't need in 2D. It hence is assumed to
 // implicitly be [0, 0, 1]
-uniform mat3x2 modelview_matrix;
+layout (location = 0) uniform mat3x2 modelview_matrix;
 
 void main() {
     // Multiply input position by the rotscale part of the matrix and then manually translate by
@@ -47,34 +49,29 @@ void main() {
 }
 )";
 
-static const char fragment_shader[] = R"(
-#version 150 core
+constexpr char fragment_shader[] = R"(
+#version 430 core
 
-in vec2 frag_tex_coord;
-out vec4 color;
+layout (location = 0) in vec2 frag_tex_coord;
+layout (location = 0) out vec4 color;
 
-uniform sampler2D color_texture;
+layout (binding = 0) uniform sampler2D color_texture;
 
 void main() {
-    // Swap RGBA -> ABGR so we don't have to do this on the CPU. This needs to change if we have to
-    // support more framebuffer pixel formats.
     color = texture(color_texture, frag_tex_coord);
 }
 )";
 
-/**
- * Vertex structure that the drawn screen rectangles are composed of.
- */
-struct ScreenRectVertex {
-    ScreenRectVertex(GLfloat x, GLfloat y, GLfloat u, GLfloat v) {
-        position[0] = x;
-        position[1] = y;
-        tex_coord[0] = u;
-        tex_coord[1] = v;
-    }
+constexpr GLint PositionLocation = 0;
+constexpr GLint TexCoordLocation = 1;
+constexpr GLint ModelViewMatrixLocation = 0;
 
-    GLfloat position[2];
-    GLfloat tex_coord[2];
+struct ScreenRectVertex {
+    constexpr ScreenRectVertex(GLfloat x, GLfloat y, GLfloat u, GLfloat v)
+        : position{{x, y}}, tex_coord{{u, v}} {}
+
+    std::array<GLfloat, 2> position;
+    std::array<GLfloat, 2> tex_coord;
 };
 
 /**
@@ -84,18 +81,82 @@ struct ScreenRectVertex {
  * The projection part of the matrix is trivial, hence these operations are represented
  * by a 3x2 matrix.
  */
-static std::array<GLfloat, 3 * 2> MakeOrthographicMatrix(const float width, const float height) {
+std::array<GLfloat, 3 * 2> MakeOrthographicMatrix(float width, float height) {
     std::array<GLfloat, 3 * 2> matrix; // Laid out in column-major order
 
     // clang-format off
-    matrix[0] = 2.f / width; matrix[2] = 0.f;           matrix[4] = -1.f;
-    matrix[1] = 0.f;         matrix[3] = -2.f / height; matrix[5] = 1.f;
+    matrix[0] = 2.f / width; matrix[2] =  0.f;          matrix[4] = -1.f;
+    matrix[1] = 0.f;         matrix[3] = -2.f / height; matrix[5] =  1.f;
     // Last matrix row is implicitly assumed to be [0, 0, 1].
     // clang-format on
 
     return matrix;
 }
 
+const char* GetSource(GLenum source) {
+    switch (source) {
+    case GL_DEBUG_SOURCE_API:
+        return "API";
+    case GL_DEBUG_SOURCE_WINDOW_SYSTEM:
+        return "WINDOW_SYSTEM";
+    case GL_DEBUG_SOURCE_SHADER_COMPILER:
+        return "SHADER_COMPILER";
+    case GL_DEBUG_SOURCE_THIRD_PARTY:
+        return "THIRD_PARTY";
+    case GL_DEBUG_SOURCE_APPLICATION:
+        return "APPLICATION";
+    case GL_DEBUG_SOURCE_OTHER:
+        return "OTHER";
+    default:
+        UNREACHABLE();
+        return "Unknown source";
+    }
+}
+
+const char* GetType(GLenum type) {
+    switch (type) {
+    case GL_DEBUG_TYPE_ERROR:
+        return "ERROR";
+    case GL_DEBUG_TYPE_DEPRECATED_BEHAVIOR:
+        return "DEPRECATED_BEHAVIOR";
+    case GL_DEBUG_TYPE_UNDEFINED_BEHAVIOR:
+        return "UNDEFINED_BEHAVIOR";
+    case GL_DEBUG_TYPE_PORTABILITY:
+        return "PORTABILITY";
+    case GL_DEBUG_TYPE_PERFORMANCE:
+        return "PERFORMANCE";
+    case GL_DEBUG_TYPE_OTHER:
+        return "OTHER";
+    case GL_DEBUG_TYPE_MARKER:
+        return "MARKER";
+    default:
+        UNREACHABLE();
+        return "Unknown type";
+    }
+}
+
+void APIENTRY DebugHandler(GLenum source, GLenum type, GLuint id, GLenum severity, GLsizei length,
+                           const GLchar* message, const void* user_param) {
+    const char format[] = "{} {} {}: {}";
+    const char* const str_source = GetSource(source);
+    const char* const str_type = GetType(type);
+
+    switch (severity) {
+    case GL_DEBUG_SEVERITY_HIGH:
+        LOG_CRITICAL(Render_OpenGL, format, str_source, str_type, id, message);
+        break;
+    case GL_DEBUG_SEVERITY_MEDIUM:
+        LOG_WARNING(Render_OpenGL, format, str_source, str_type, id, message);
+        break;
+    case GL_DEBUG_SEVERITY_NOTIFICATION:
+    case GL_DEBUG_SEVERITY_LOW:
+        LOG_DEBUG(Render_OpenGL, format, str_source, str_type, id, message);
+        break;
+    }
+}
+
+} // Anonymous namespace
+
 RendererOpenGL::RendererOpenGL(Core::Frontend::EmuWindow& emu_window, Core::System& system)
     : VideoCore::RendererBase{emu_window}, emu_window{emu_window}, system{system} {}
 
@@ -138,9 +199,6 @@ void RendererOpenGL::SwapBuffers(const Tegra::FramebufferConfig* framebuffer) {
     prev_state.Apply();
 }
 
-/**
- * Loads framebuffer from emulated memory into the active OpenGL texture.
- */
 void RendererOpenGL::LoadFBToScreenInfo(const Tegra::FramebufferConfig& framebuffer) {
     // Framebuffer orientation handling
     framebuffer_transform_flags = framebuffer.transform_flags;
@@ -181,19 +239,12 @@ void RendererOpenGL::LoadFBToScreenInfo(const Tegra::FramebufferConfig& framebuf
     glPixelStorei(GL_UNPACK_ROW_LENGTH, 0);
 }
 
-/**
- * Fills active OpenGL texture with the given RGB color. Since the color is solid, the texture can
- * be 1x1 but will stretch across whatever it's rendered on.
- */
 void RendererOpenGL::LoadColorToActiveGLTexture(u8 color_r, u8 color_g, u8 color_b, u8 color_a,
                                                 const TextureInfo& texture) {
     const u8 framebuffer_data[4] = {color_a, color_b, color_g, color_r};
     glClearTexImage(texture.resource.handle, 0, GL_RGBA, GL_UNSIGNED_BYTE, framebuffer_data);
 }
 
-/**
- * Initializes the OpenGL state and creates persistent objects.
- */
 void RendererOpenGL::InitOpenGLObjects() {
     glClearColor(Settings::values.bg_red, Settings::values.bg_green, Settings::values.bg_blue,
                  0.0f);
@@ -203,10 +254,6 @@ void RendererOpenGL::InitOpenGLObjects() {
     state.draw.shader_program = shader.handle;
     state.AllDirty();
     state.Apply();
-    uniform_modelview_matrix = glGetUniformLocation(shader.handle, "modelview_matrix");
-    uniform_color_texture = glGetUniformLocation(shader.handle, "color_texture");
-    attrib_position = glGetAttribLocation(shader.handle, "vert_position");
-    attrib_tex_coord = glGetAttribLocation(shader.handle, "vert_tex_coord");
 
     // Generate VBO handle for drawing
     vertex_buffer.Create();
@@ -217,14 +264,14 @@ void RendererOpenGL::InitOpenGLObjects() {
 
     // Attach vertex data to VAO
     glNamedBufferData(vertex_buffer.handle, sizeof(ScreenRectVertex) * 4, nullptr, GL_STREAM_DRAW);
-    glVertexArrayAttribFormat(vertex_array.handle, attrib_position, 2, GL_FLOAT, GL_FALSE,
+    glVertexArrayAttribFormat(vertex_array.handle, PositionLocation, 2, GL_FLOAT, GL_FALSE,
                               offsetof(ScreenRectVertex, position));
-    glVertexArrayAttribFormat(vertex_array.handle, attrib_tex_coord, 2, GL_FLOAT, GL_FALSE,
+    glVertexArrayAttribFormat(vertex_array.handle, TexCoordLocation, 2, GL_FLOAT, GL_FALSE,
                               offsetof(ScreenRectVertex, tex_coord));
-    glVertexArrayAttribBinding(vertex_array.handle, attrib_position, 0);
-    glVertexArrayAttribBinding(vertex_array.handle, attrib_tex_coord, 0);
-    glEnableVertexArrayAttrib(vertex_array.handle, attrib_position);
-    glEnableVertexArrayAttrib(vertex_array.handle, attrib_tex_coord);
+    glVertexArrayAttribBinding(vertex_array.handle, PositionLocation, 0);
+    glVertexArrayAttribBinding(vertex_array.handle, TexCoordLocation, 0);
+    glEnableVertexArrayAttrib(vertex_array.handle, PositionLocation);
+    glEnableVertexArrayAttrib(vertex_array.handle, TexCoordLocation);
     glVertexArrayVertexBuffer(vertex_array.handle, 0, vertex_buffer.handle, 0,
                               sizeof(ScreenRectVertex));
 
@@ -331,18 +378,18 @@ void RendererOpenGL::DrawScreenTriangles(const ScreenInfo& screen_info, float x,
                   static_cast<f32>(screen_info.texture.height);
     }
 
-    std::array<ScreenRectVertex, 4> vertices = {{
+    const std::array vertices = {
         ScreenRectVertex(x, y, texcoords.top * scale_u, left * scale_v),
         ScreenRectVertex(x + w, y, texcoords.bottom * scale_u, left * scale_v),
         ScreenRectVertex(x, y + h, texcoords.top * scale_u, right * scale_v),
         ScreenRectVertex(x + w, y + h, texcoords.bottom * scale_u, right * scale_v),
-    }};
+    };
 
     state.textures[0] = screen_info.display_texture;
     state.framebuffer_srgb.enabled = screen_info.display_srgb;
     state.AllDirty();
     state.Apply();
-    glNamedBufferSubData(vertex_buffer.handle, 0, sizeof(vertices), vertices.data());
+    glNamedBufferSubData(vertex_buffer.handle, 0, sizeof(vertices), std::data(vertices));
     glDrawArrays(GL_TRIANGLE_STRIP, 0, 4);
     // Restore default state
     state.framebuffer_srgb.enabled = false;
@@ -351,9 +398,6 @@ void RendererOpenGL::DrawScreenTriangles(const ScreenInfo& screen_info, float x,
     state.Apply();
 }
 
-/**
- * Draws the emulated screens to the emulator window.
- */
 void RendererOpenGL::DrawScreen(const Layout::FramebufferLayout& layout) {
     if (renderer_settings.set_background_color) {
         // Update background color before drawing
@@ -367,21 +411,17 @@ void RendererOpenGL::DrawScreen(const Layout::FramebufferLayout& layout) {
     glClear(GL_COLOR_BUFFER_BIT);
 
     // Set projection matrix
-    std::array<GLfloat, 3 * 2> ortho_matrix =
-        MakeOrthographicMatrix((float)layout.width, (float)layout.height);
-    glUniformMatrix3x2fv(uniform_modelview_matrix, 1, GL_FALSE, ortho_matrix.data());
+    const std::array ortho_matrix =
+        MakeOrthographicMatrix(static_cast<float>(layout.width), static_cast<float>(layout.height));
+    glUniformMatrix3x2fv(ModelViewMatrixLocation, 1, GL_FALSE, ortho_matrix.data());
 
-    // Bind texture in Texture Unit 0
-    glActiveTexture(GL_TEXTURE0);
-    glUniform1i(uniform_color_texture, 0);
-
-    DrawScreenTriangles(screen_info, (float)screen.left, (float)screen.top,
-                        (float)screen.GetWidth(), (float)screen.GetHeight());
+    DrawScreenTriangles(screen_info, static_cast<float>(screen.left),
+                        static_cast<float>(screen.top), static_cast<float>(screen.GetWidth()),
+                        static_cast<float>(screen.GetHeight()));
 
     m_current_frame++;
 }
 
-/// Updates the framerate
 void RendererOpenGL::UpdateFramerate() {}
 
 void RendererOpenGL::CaptureScreenshot() {
@@ -418,63 +458,6 @@ void RendererOpenGL::CaptureScreenshot() {
     renderer_settings.screenshot_requested = false;
 }
 
-static const char* GetSource(GLenum source) {
-#define RET(s)                                                                                     \
-    case GL_DEBUG_SOURCE_##s:                                                                      \
-        return #s
-    switch (source) {
-        RET(API);
-        RET(WINDOW_SYSTEM);
-        RET(SHADER_COMPILER);
-        RET(THIRD_PARTY);
-        RET(APPLICATION);
-        RET(OTHER);
-    default:
-        UNREACHABLE();
-        return "Unknown source";
-    }
-#undef RET
-}
-
-static const char* GetType(GLenum type) {
-#define RET(t)                                                                                     \
-    case GL_DEBUG_TYPE_##t:                                                                        \
-        return #t
-    switch (type) {
-        RET(ERROR);
-        RET(DEPRECATED_BEHAVIOR);
-        RET(UNDEFINED_BEHAVIOR);
-        RET(PORTABILITY);
-        RET(PERFORMANCE);
-        RET(OTHER);
-        RET(MARKER);
-    default:
-        UNREACHABLE();
-        return "Unknown type";
-    }
-#undef RET
-}
-
-static void APIENTRY DebugHandler(GLenum source, GLenum type, GLuint id, GLenum severity,
-                                  GLsizei length, const GLchar* message, const void* user_param) {
-    const char format[] = "{} {} {}: {}";
-    const char* const str_source = GetSource(source);
-    const char* const str_type = GetType(type);
-
-    switch (severity) {
-    case GL_DEBUG_SEVERITY_HIGH:
-        LOG_CRITICAL(Render_OpenGL, format, str_source, str_type, id, message);
-        break;
-    case GL_DEBUG_SEVERITY_MEDIUM:
-        LOG_WARNING(Render_OpenGL, format, str_source, str_type, id, message);
-        break;
-    case GL_DEBUG_SEVERITY_NOTIFICATION:
-    case GL_DEBUG_SEVERITY_LOW:
-        LOG_DEBUG(Render_OpenGL, format, str_source, str_type, id, message);
-        break;
-    }
-}
-
 bool RendererOpenGL::Init() {
     Core::Frontend::ScopeAcquireWindowContext acquire_context{render_window};
 
@@ -495,7 +478,6 @@ bool RendererOpenGL::Init() {
     return true;
 }
 
-/// Shutdown the renderer
 void RendererOpenGL::ShutDown() {}
 
 } // namespace OpenGL

src/video_core/renderer_opengl/renderer_opengl.h

@@ -59,21 +59,31 @@ public:
     void ShutDown() override;
 
 private:
+    /// Initializes the OpenGL state and creates persistent objects.
     void InitOpenGLObjects();
+
     void AddTelemetryFields();
+
     void CreateRasterizer();
 
     void ConfigureFramebufferTexture(TextureInfo& texture,
                                      const Tegra::FramebufferConfig& framebuffer);
+
+    /// Draws the emulated screens to the emulator window.
     void DrawScreen(const Layout::FramebufferLayout& layout);
+
     void DrawScreenTriangles(const ScreenInfo& screen_info, float x, float y, float w, float h);
+
+    /// Updates the framerate.
     void UpdateFramerate();
 
     void CaptureScreenshot();
 
-    // Loads framebuffer from emulated memory into the display information structure
+    /// Loads framebuffer from emulated memory into the active OpenGL texture.
     void LoadFBToScreenInfo(const Tegra::FramebufferConfig& framebuffer);
-    // Fills active OpenGL texture with the given RGBA color.
+
+    /// Fills active OpenGL texture with the given RGB color.Since the color is solid, the texture
+    /// can be 1x1 but will stretch across whatever it's rendered on.
     void LoadColorToActiveGLTexture(u8 color_r, u8 color_g, u8 color_b, u8 color_a,
                                     const TextureInfo& texture);
 
@@ -94,14 +104,6 @@ private:
     /// OpenGL framebuffer data
     std::vector<u8> gl_framebuffer_data;
 
-    // Shader uniform location indices
-    GLuint uniform_modelview_matrix;
-    GLuint uniform_color_texture;
-
-    // Shader attribute input indices
-    GLuint attrib_position;
-    GLuint attrib_tex_coord;
-
     /// Used for transforming the framebuffer orientation
     Tegra::FramebufferConfig::TransformFlags framebuffer_transform_flags;
     Common::Rectangle<int> framebuffer_crop_rect;

src/video_core/renderer_vulkan/declarations.h

@@ -4,6 +4,17 @@
 
 #pragma once
 
+namespace vk {
+class DispatchLoaderDynamic;
+}
+
+namespace Vulkan {
+constexpr vk::DispatchLoaderDynamic* dont_use_me_dld = nullptr;
+}
+
+#define VULKAN_HPP_DEFAULT_DISPATCHER (*::Vulkan::dont_use_me_dld)
+#define VULKAN_HPP_ENABLE_DYNAMIC_LOADER_TOOL 0
+#define VULKAN_HPP_DISPATCH_LOADER_DYNAMIC 1
 #include <vulkan/vulkan.hpp>
 
 namespace Vulkan {
@@ -41,5 +52,7 @@ using UniqueSemaphore = UniqueHandle<vk::Semaphore>;
 using UniqueShaderModule = UniqueHandle<vk::ShaderModule>;
 using UniqueSwapchainKHR = UniqueHandle<vk::SwapchainKHR>;
 using UniqueValidationCacheEXT = UniqueHandle<vk::ValidationCacheEXT>;
+using UniqueDebugReportCallbackEXT = UniqueHandle<vk::DebugReportCallbackEXT>;
+using UniqueDebugUtilsMessengerEXT = UniqueHandle<vk::DebugUtilsMessengerEXT>;
 
 } // namespace Vulkan

src/video_core/renderer_vulkan/maxwell_to_vk.cpp

@@ -44,7 +44,8 @@ vk::SamplerMipmapMode MipmapMode(Tegra::Texture::TextureMipmapFilter mipmap_filt
     return {};
 }
 
-vk::SamplerAddressMode WrapMode(Tegra::Texture::WrapMode wrap_mode) {
+vk::SamplerAddressMode WrapMode(Tegra::Texture::WrapMode wrap_mode,
+                                Tegra::Texture::TextureFilter filter) {
     switch (wrap_mode) {
     case Tegra::Texture::WrapMode::Wrap:
         return vk::SamplerAddressMode::eRepeat;
@@ -55,10 +56,15 @@ vk::SamplerAddressMode WrapMode(Tegra::Texture::WrapMode wrap_mode) {
     case Tegra::Texture::WrapMode::Border:
         return vk::SamplerAddressMode::eClampToBorder;
     case Tegra::Texture::WrapMode::Clamp:
-        // TODO(Rodrigo): GL_CLAMP was removed as of OpenGL 3.1, to implement GL_CLAMP, we can use
-        // eClampToBorder to get the border color of the texture, and then sample the edge to
-        // manually mix them. However the shader part of this is not yet implemented.
-        return vk::SamplerAddressMode::eClampToBorder;
+        // TODO(Rodrigo): Emulate GL_CLAMP properly
+        switch (filter) {
+        case Tegra::Texture::TextureFilter::Nearest:
+            return vk::SamplerAddressMode::eClampToEdge;
+        case Tegra::Texture::TextureFilter::Linear:
+            return vk::SamplerAddressMode::eClampToBorder;
+        }
+        UNREACHABLE();
+        return vk::SamplerAddressMode::eClampToEdge;
     case Tegra::Texture::WrapMode::MirrorOnceClampToEdge:
         return vk::SamplerAddressMode::eMirrorClampToEdge;
     case Tegra::Texture::WrapMode::MirrorOnceBorder:
@@ -96,106 +102,140 @@ vk::CompareOp DepthCompareFunction(Tegra::Texture::DepthCompareFunc depth_compar
 
 } // namespace Sampler
 
+namespace {
+
+enum : u32 { Attachable = 1, Storage = 2 };
+
 struct FormatTuple {
     vk::Format format; ///< Vulkan format
-    bool attachable;   ///< True when this format can be used as an attachment
-};
-
-static constexpr std::array<FormatTuple, VideoCore::Surface::MaxPixelFormat> tex_format_tuples = {{
-    {vk::Format::eA8B8G8R8UnormPack32, true},    // ABGR8U
-    {vk::Format::eUndefined, false},             // ABGR8S
-    {vk::Format::eUndefined, false},             // ABGR8UI
-    {vk::Format::eB5G6R5UnormPack16, false},     // B5G6R5U
-    {vk::Format::eA2B10G10R10UnormPack32, true}, // A2B10G10R10U
-    {vk::Format::eUndefined, false},             // A1B5G5R5U
-    {vk::Format::eR8Unorm, true},                // R8U
-    {vk::Format::eUndefined, false},             // R8UI
-    {vk::Format::eUndefined, false},             // RGBA16F
-    {vk::Format::eUndefined, false},             // RGBA16U
-    {vk::Format::eUndefined, false},             // RGBA16UI
-    {vk::Format::eUndefined, false},             // R11FG11FB10F
-    {vk::Format::eUndefined, false},             // RGBA32UI
-    {vk::Format::eBc1RgbaUnormBlock, false},     // DXT1
-    {vk::Format::eBc2UnormBlock, false},         // DXT23
-    {vk::Format::eBc3UnormBlock, false},         // DXT45
-    {vk::Format::eBc4UnormBlock, false},         // DXN1
-    {vk::Format::eUndefined, false},             // DXN2UNORM
-    {vk::Format::eUndefined, false},             // DXN2SNORM
-    {vk::Format::eUndefined, false},             // BC7U
-    {vk::Format::eUndefined, false},             // BC6H_UF16
-    {vk::Format::eUndefined, false},             // BC6H_SF16
-    {vk::Format::eUndefined, false},             // ASTC_2D_4X4
-    {vk::Format::eUndefined, false},             // BGRA8
-    {vk::Format::eUndefined, false},             // RGBA32F
-    {vk::Format::eUndefined, false},             // RG32F
-    {vk::Format::eUndefined, false},             // R32F
-    {vk::Format::eUndefined, false},             // R16F
-    {vk::Format::eUndefined, false},             // R16U
-    {vk::Format::eUndefined, false},             // R16S
-    {vk::Format::eUndefined, false},             // R16UI
-    {vk::Format::eUndefined, false},             // R16I
-    {vk::Format::eUndefined, false},             // RG16
-    {vk::Format::eUndefined, false},             // RG16F
-    {vk::Format::eUndefined, false},             // RG16UI
-    {vk::Format::eUndefined, false},             // RG16I
-    {vk::Format::eUndefined, false},             // RG16S
-    {vk::Format::eUndefined, false},             // RGB32F
-    {vk::Format::eA8B8G8R8SrgbPack32, true},     // RGBA8_SRGB
-    {vk::Format::eUndefined, false},             // RG8U
-    {vk::Format::eUndefined, false},             // RG8S
-    {vk::Format::eUndefined, false},             // RG32UI
-    {vk::Format::eUndefined, false},             // RGBX16F
-    {vk::Format::eUndefined, false},             // R32UI
-    {vk::Format::eUndefined, false},             // ASTC_2D_8X8
-    {vk::Format::eUndefined, false},             // ASTC_2D_8X5
-    {vk::Format::eUndefined, false},             // ASTC_2D_5X4
-
-    // Compressed sRGB formats
-    {vk::Format::eUndefined, false}, // BGRA8_SRGB
-    {vk::Format::eUndefined, false}, // DXT1_SRGB
-    {vk::Format::eUndefined, false}, // DXT23_SRGB
-    {vk::Format::eUndefined, false}, // DXT45_SRGB
-    {vk::Format::eUndefined, false}, // BC7U_SRGB
-    {vk::Format::eUndefined, false}, // ASTC_2D_4X4_SRGB
-    {vk::Format::eUndefined, false}, // ASTC_2D_8X8_SRGB
-    {vk::Format::eUndefined, false}, // ASTC_2D_8X5_SRGB
-    {vk::Format::eUndefined, false}, // ASTC_2D_5X4_SRGB
-    {vk::Format::eUndefined, false}, // ASTC_2D_5X5
-    {vk::Format::eUndefined, false}, // ASTC_2D_5X5_SRGB
-    {vk::Format::eUndefined, false}, // ASTC_2D_10X8
-    {vk::Format::eUndefined, false}, // ASTC_2D_10X8_SRGB
+    int usage;         ///< Describes image format usage
+} constexpr tex_format_tuples[] = {
+    {vk::Format::eA8B8G8R8UnormPack32, Attachable | Storage},    // ABGR8U
+    {vk::Format::eA8B8G8R8SnormPack32, Attachable | Storage},    // ABGR8S
+    {vk::Format::eA8B8G8R8UintPack32, Attachable | Storage},     // ABGR8UI
+    {vk::Format::eB5G6R5UnormPack16, {}},                        // B5G6R5U
+    {vk::Format::eA2B10G10R10UnormPack32, Attachable | Storage}, // A2B10G10R10U
+    {vk::Format::eA1R5G5B5UnormPack16, Attachable | Storage},    // A1B5G5R5U (flipped with swizzle)
+    {vk::Format::eR8Unorm, Attachable | Storage},                // R8U
+    {vk::Format::eR8Uint, Attachable | Storage},                 // R8UI
+    {vk::Format::eR16G16B16A16Sfloat, Attachable | Storage},     // RGBA16F
+    {vk::Format::eR16G16B16A16Unorm, Attachable | Storage},      // RGBA16U
+    {vk::Format::eR16G16B16A16Uint, Attachable | Storage},       // RGBA16UI
+    {vk::Format::eB10G11R11UfloatPack32, Attachable | Storage},  // R11FG11FB10F
+    {vk::Format::eR32G32B32A32Uint, Attachable | Storage},       // RGBA32UI
+    {vk::Format::eBc1RgbaUnormBlock, {}},                        // DXT1
+    {vk::Format::eBc2UnormBlock, {}},                            // DXT23
+    {vk::Format::eBc3UnormBlock, {}},                            // DXT45
+    {vk::Format::eBc4UnormBlock, {}},                            // DXN1
+    {vk::Format::eBc5UnormBlock, {}},                            // DXN2UNORM
+    {vk::Format::eBc5SnormBlock, {}},                            // DXN2SNORM
+    {vk::Format::eBc7UnormBlock, {}},                            // BC7U
+    {vk::Format::eBc6HUfloatBlock, {}},                          // BC6H_UF16
+    {vk::Format::eBc6HSfloatBlock, {}},                          // BC6H_SF16
+    {vk::Format::eAstc4x4UnormBlock, {}},                        // ASTC_2D_4X4
+    {vk::Format::eB8G8R8A8Unorm, {}},                            // BGRA8
+    {vk::Format::eR32G32B32A32Sfloat, Attachable | Storage},     // RGBA32F
+    {vk::Format::eR32G32Sfloat, Attachable | Storage},           // RG32F
+    {vk::Format::eR32Sfloat, Attachable | Storage},              // R32F
+    {vk::Format::eR16Sfloat, Attachable | Storage},              // R16F
+    {vk::Format::eR16Unorm, Attachable | Storage},               // R16U
+    {vk::Format::eUndefined, {}},                                // R16S
+    {vk::Format::eUndefined, {}},                                // R16UI
+    {vk::Format::eUndefined, {}},                                // R16I
+    {vk::Format::eR16G16Unorm, Attachable | Storage},            // RG16
+    {vk::Format::eR16G16Sfloat, Attachable | Storage},           // RG16F
+    {vk::Format::eUndefined, {}},                                // RG16UI
+    {vk::Format::eUndefined, {}},                                // RG16I
+    {vk::Format::eR16G16Snorm, Attachable | Storage},            // RG16S
+    {vk::Format::eUndefined, {}},                                // RGB32F
+    {vk::Format::eR8G8B8A8Srgb, Attachable},                     // RGBA8_SRGB
+    {vk::Format::eR8G8Unorm, Attachable | Storage},              // RG8U
+    {vk::Format::eR8G8Snorm, Attachable | Storage},              // RG8S
+    {vk::Format::eR32G32Uint, Attachable | Storage},             // RG32UI
+    {vk::Format::eUndefined, {}},                                // RGBX16F
+    {vk::Format::eR32Uint, Attachable | Storage},                // R32UI
+    {vk::Format::eAstc8x8UnormBlock, {}},                        // ASTC_2D_8X8
+    {vk::Format::eUndefined, {}},                                // ASTC_2D_8X5
+    {vk::Format::eUndefined, {}},                                // ASTC_2D_5X4
+    {vk::Format::eUndefined, {}},                                // BGRA8_SRGB
+    {vk::Format::eBc1RgbaSrgbBlock, {}},                         // DXT1_SRGB
+    {vk::Format::eUndefined, {}},                                // DXT23_SRGB
+    {vk::Format::eBc3SrgbBlock, {}},                             // DXT45_SRGB
+    {vk::Format::eBc7SrgbBlock, {}},                             // BC7U_SRGB
+    {vk::Format::eR4G4B4A4UnormPack16, Attachable},              // R4G4B4A4U
+    {vk::Format::eAstc4x4SrgbBlock, {}},                         // ASTC_2D_4X4_SRGB
+    {vk::Format::eAstc8x8SrgbBlock, {}},                         // ASTC_2D_8X8_SRGB
+    {vk::Format::eAstc8x5SrgbBlock, {}},                         // ASTC_2D_8X5_SRGB
+    {vk::Format::eAstc5x4SrgbBlock, {}},                         // ASTC_2D_5X4_SRGB
+    {vk::Format::eAstc5x5UnormBlock, {}},                        // ASTC_2D_5X5
+    {vk::Format::eAstc5x5SrgbBlock, {}},                         // ASTC_2D_5X5_SRGB
+    {vk::Format::eAstc10x8UnormBlock, {}},                       // ASTC_2D_10X8
+    {vk::Format::eAstc10x8SrgbBlock, {}},                        // ASTC_2D_10X8_SRGB
+    {vk::Format::eAstc6x6UnormBlock, {}},                        // ASTC_2D_6X6
+    {vk::Format::eAstc6x6SrgbBlock, {}},                         // ASTC_2D_6X6_SRGB
+    {vk::Format::eAstc10x10UnormBlock, {}},                      // ASTC_2D_10X10
+    {vk::Format::eAstc10x10SrgbBlock, {}},                       // ASTC_2D_10X10_SRGB
+    {vk::Format::eAstc12x12UnormBlock, {}},                      // ASTC_2D_12X12
+    {vk::Format::eAstc12x12SrgbBlock, {}},                       // ASTC_2D_12X12_SRGB
+    {vk::Format::eAstc8x6UnormBlock, {}},                        // ASTC_2D_8X6
+    {vk::Format::eAstc8x6SrgbBlock, {}},                         // ASTC_2D_8X6_SRGB
+    {vk::Format::eAstc6x5UnormBlock, {}},                        // ASTC_2D_6X5
+    {vk::Format::eAstc6x5SrgbBlock, {}},                         // ASTC_2D_6X5_SRGB
+    {vk::Format::eE5B9G9R9UfloatPack32, {}},                     // E5B9G9R9F
 
     // Depth formats
-    {vk::Format::eD32Sfloat, true}, // Z32F
-    {vk::Format::eD16Unorm, true},  // Z16
+    {vk::Format::eD32Sfloat, Attachable}, // Z32F
+    {vk::Format::eD16Unorm, Attachable},  // Z16
 
     // DepthStencil formats
-    {vk::Format::eD24UnormS8Uint, true}, // Z24S8
-    {vk::Format::eD24UnormS8Uint, true}, // S8Z24 (emulated)
-    {vk::Format::eUndefined, false},     // Z32FS8
-}};
+    {vk::Format::eD24UnormS8Uint, Attachable},  // Z24S8
+    {vk::Format::eD24UnormS8Uint, Attachable},  // S8Z24 (emulated)
+    {vk::Format::eD32SfloatS8Uint, Attachable}, // Z32FS8
+};
+static_assert(std::size(tex_format_tuples) == VideoCore::Surface::MaxPixelFormat);
 
-static constexpr bool IsZetaFormat(PixelFormat pixel_format) {
+constexpr bool IsZetaFormat(PixelFormat pixel_format) {
     return pixel_format >= PixelFormat::MaxColorFormat &&
            pixel_format < PixelFormat::MaxDepthStencilFormat;
 }
 
-std::pair<vk::Format, bool> SurfaceFormat(const VKDevice& device, FormatType format_type,
-                                          PixelFormat pixel_format) {
-    ASSERT(static_cast<std::size_t>(pixel_format) < tex_format_tuples.size());
+} // Anonymous namespace
 
-    const auto tuple = tex_format_tuples[static_cast<u32>(pixel_format)];
-    UNIMPLEMENTED_IF_MSG(tuple.format == vk::Format::eUndefined,
-                         "Unimplemented texture format with pixel format={}",
-                         static_cast<u32>(pixel_format));
+FormatInfo SurfaceFormat(const VKDevice& device, FormatType format_type, PixelFormat pixel_format) {
+    ASSERT(static_cast<std::size_t>(pixel_format) < std::size(tex_format_tuples));
 
-    auto usage = vk::FormatFeatureFlagBits::eSampledImage |
-                 vk::FormatFeatureFlagBits::eTransferDst | vk::FormatFeatureFlagBits::eTransferSrc;
-    if (tuple.attachable) {
-        usage |= IsZetaFormat(pixel_format) ? vk::FormatFeatureFlagBits::eDepthStencilAttachment
-                                            : vk::FormatFeatureFlagBits::eColorAttachment;
+    auto tuple = tex_format_tuples[static_cast<std::size_t>(pixel_format)];
+    if (tuple.format == vk::Format::eUndefined) {
+        UNIMPLEMENTED_MSG("Unimplemented texture format with pixel format={}",
+                          static_cast<u32>(pixel_format));
+        return {vk::Format::eA8B8G8R8UnormPack32, true, true};
     }
-    return {device.GetSupportedFormat(tuple.format, usage, format_type), tuple.attachable};
+
+    // Use ABGR8 on hardware that doesn't support ASTC natively
+    if (!device.IsOptimalAstcSupported() && VideoCore::Surface::IsPixelFormatASTC(pixel_format)) {
+        tuple.format = VideoCore::Surface::IsPixelFormatSRGB(pixel_format)
+                           ? vk::Format::eA8B8G8R8SrgbPack32
+                           : vk::Format::eA8B8G8R8UnormPack32;
+    }
+    const bool attachable = tuple.usage & Attachable;
+    const bool storage = tuple.usage & Storage;
+
+    vk::FormatFeatureFlags usage;
+    if (format_type == FormatType::Buffer) {
+        usage = vk::FormatFeatureFlagBits::eStorageTexelBuffer |
+                vk::FormatFeatureFlagBits::eUniformTexelBuffer;
+    } else {
+        usage = vk::FormatFeatureFlagBits::eSampledImage | vk::FormatFeatureFlagBits::eTransferDst |
+                vk::FormatFeatureFlagBits::eTransferSrc;
+        if (attachable) {
+            usage |= IsZetaFormat(pixel_format) ? vk::FormatFeatureFlagBits::eDepthStencilAttachment
+                                                : vk::FormatFeatureFlagBits::eColorAttachment;
+        }
+        if (storage) {
+            usage |= vk::FormatFeatureFlagBits::eStorageImage;
+        }
+    }
+    return {device.GetSupportedFormat(tuple.format, usage, format_type), attachable, storage};
 }
 
 vk::ShaderStageFlagBits ShaderStage(Tegra::Engines::ShaderType stage) {
@@ -215,7 +255,8 @@ vk::ShaderStageFlagBits ShaderStage(Tegra::Engines::ShaderType stage) {
     return {};
 }
 
-vk::PrimitiveTopology PrimitiveTopology(Maxwell::PrimitiveTopology topology) {
+vk::PrimitiveTopology PrimitiveTopology([[maybe_unused]] const VKDevice& device,
+                                        Maxwell::PrimitiveTopology topology) {
     switch (topology) {
     case Maxwell::PrimitiveTopology::Points:
         return vk::PrimitiveTopology::ePointList;
@@ -227,6 +268,13 @@ vk::PrimitiveTopology PrimitiveTopology(Maxwell::PrimitiveTopology topology) {
         return vk::PrimitiveTopology::eTriangleList;
     case Maxwell::PrimitiveTopology::TriangleStrip:
         return vk::PrimitiveTopology::eTriangleStrip;
+    case Maxwell::PrimitiveTopology::TriangleFan:
+        return vk::PrimitiveTopology::eTriangleFan;
+    case Maxwell::PrimitiveTopology::Quads:
+        // TODO(Rodrigo): Use VK_PRIMITIVE_TOPOLOGY_QUAD_LIST_EXT whenever it releases
+        return vk::PrimitiveTopology::eTriangleList;
+    case Maxwell::PrimitiveTopology::Patches:
+        return vk::PrimitiveTopology::ePatchList;
     default:
         UNIMPLEMENTED_MSG("Unimplemented topology={}", static_cast<u32>(topology));
         return {};
@@ -236,37 +284,111 @@ vk::PrimitiveTopology PrimitiveTopology(Maxwell::PrimitiveTopology topology) {
 vk::Format VertexFormat(Maxwell::VertexAttribute::Type type, Maxwell::VertexAttribute::Size size) {
     switch (type) {
     case Maxwell::VertexAttribute::Type::SignedNorm:
+        switch (size) {
+        case Maxwell::VertexAttribute::Size::Size_8:
+            return vk::Format::eR8Snorm;
+        case Maxwell::VertexAttribute::Size::Size_8_8:
+            return vk::Format::eR8G8Snorm;
+        case Maxwell::VertexAttribute::Size::Size_8_8_8:
+            return vk::Format::eR8G8B8Snorm;
+        case Maxwell::VertexAttribute::Size::Size_8_8_8_8:
+            return vk::Format::eR8G8B8A8Snorm;
+        case Maxwell::VertexAttribute::Size::Size_16:
+            return vk::Format::eR16Snorm;
+        case Maxwell::VertexAttribute::Size::Size_16_16:
+            return vk::Format::eR16G16Snorm;
+        case Maxwell::VertexAttribute::Size::Size_16_16_16:
+            return vk::Format::eR16G16B16Snorm;
+        case Maxwell::VertexAttribute::Size::Size_16_16_16_16:
+            return vk::Format::eR16G16B16A16Snorm;
+        case Maxwell::VertexAttribute::Size::Size_10_10_10_2:
+            return vk::Format::eA2B10G10R10SnormPack32;
+        default:
+            break;
+        }
         break;
     case Maxwell::VertexAttribute::Type::UnsignedNorm:
         switch (size) {
+        case Maxwell::VertexAttribute::Size::Size_8:
+            return vk::Format::eR8Unorm;
+        case Maxwell::VertexAttribute::Size::Size_8_8:
+            return vk::Format::eR8G8Unorm;
+        case Maxwell::VertexAttribute::Size::Size_8_8_8:
+            return vk::Format::eR8G8B8Unorm;
         case Maxwell::VertexAttribute::Size::Size_8_8_8_8:
             return vk::Format::eR8G8B8A8Unorm;
+        case Maxwell::VertexAttribute::Size::Size_16:
+            return vk::Format::eR16Unorm;
+        case Maxwell::VertexAttribute::Size::Size_16_16:
+            return vk::Format::eR16G16Unorm;
+        case Maxwell::VertexAttribute::Size::Size_16_16_16:
+            return vk::Format::eR16G16B16Unorm;
+        case Maxwell::VertexAttribute::Size::Size_16_16_16_16:
+            return vk::Format::eR16G16B16A16Unorm;
         default:
             break;
         }
         break;
     case Maxwell::VertexAttribute::Type::SignedInt:
-        break;
+        switch (size) {
+        case Maxwell::VertexAttribute::Size::Size_16_16_16_16:
+            return vk::Format::eR16G16B16A16Sint;
+        case Maxwell::VertexAttribute::Size::Size_8:
+            return vk::Format::eR8Sint;
+        case Maxwell::VertexAttribute::Size::Size_8_8:
+            return vk::Format::eR8G8Sint;
+        case Maxwell::VertexAttribute::Size::Size_8_8_8:
+            return vk::Format::eR8G8B8Sint;
+        case Maxwell::VertexAttribute::Size::Size_8_8_8_8:
+            return vk::Format::eR8G8B8A8Sint;
+        case Maxwell::VertexAttribute::Size::Size_32:
+            return vk::Format::eR32Sint;
+        default:
+            break;
+        }
     case Maxwell::VertexAttribute::Type::UnsignedInt:
         switch (size) {
+        case Maxwell::VertexAttribute::Size::Size_8:
+            return vk::Format::eR8Uint;
+        case Maxwell::VertexAttribute::Size::Size_8_8:
+            return vk::Format::eR8G8Uint;
+        case Maxwell::VertexAttribute::Size::Size_8_8_8:
+            return vk::Format::eR8G8B8Uint;
+        case Maxwell::VertexAttribute::Size::Size_8_8_8_8:
+            return vk::Format::eR8G8B8A8Uint;
         case Maxwell::VertexAttribute::Size::Size_32:
             return vk::Format::eR32Uint;
         default:
             break;
         }
     case Maxwell::VertexAttribute::Type::UnsignedScaled:
+        switch (size) {
+        case Maxwell::VertexAttribute::Size::Size_8_8:
+            return vk::Format::eR8G8Uscaled;
+        default:
+            break;
+        }
+        break;
     case Maxwell::VertexAttribute::Type::SignedScaled:
         break;
     case Maxwell::VertexAttribute::Type::Float:
         switch (size) {
-        case Maxwell::VertexAttribute::Size::Size_32_32_32_32:
-            return vk::Format::eR32G32B32A32Sfloat;
-        case Maxwell::VertexAttribute::Size::Size_32_32_32:
-            return vk::Format::eR32G32B32Sfloat;
-        case Maxwell::VertexAttribute::Size::Size_32_32:
-            return vk::Format::eR32G32Sfloat;
         case Maxwell::VertexAttribute::Size::Size_32:
             return vk::Format::eR32Sfloat;
+        case Maxwell::VertexAttribute::Size::Size_32_32:
+            return vk::Format::eR32G32Sfloat;
+        case Maxwell::VertexAttribute::Size::Size_32_32_32:
+            return vk::Format::eR32G32B32Sfloat;
+        case Maxwell::VertexAttribute::Size::Size_32_32_32_32:
+            return vk::Format::eR32G32B32A32Sfloat;
+        case Maxwell::VertexAttribute::Size::Size_16:
+            return vk::Format::eR16Sfloat;
+        case Maxwell::VertexAttribute::Size::Size_16_16:
+            return vk::Format::eR16G16Sfloat;
+        case Maxwell::VertexAttribute::Size::Size_16_16_16:
+            return vk::Format::eR16G16B16Sfloat;
+        case Maxwell::VertexAttribute::Size::Size_16_16_16_16:
+            return vk::Format::eR16G16B16A16Sfloat;
         default:
             break;
         }
@@ -308,11 +430,14 @@ vk::CompareOp ComparisonOp(Maxwell::ComparisonOp comparison) {
     return {};
 }
 
-vk::IndexType IndexFormat(Maxwell::IndexFormat index_format) {
+vk::IndexType IndexFormat(const VKDevice& device, Maxwell::IndexFormat index_format) {
     switch (index_format) {
     case Maxwell::IndexFormat::UnsignedByte:
-        UNIMPLEMENTED_MSG("Vulkan does not support native u8 index format");
-        return vk::IndexType::eUint16;
+        if (!device.IsExtIndexTypeUint8Supported()) {
+            UNIMPLEMENTED_MSG("Native uint8 indices are not supported on this device");
+            return vk::IndexType::eUint16;
+        }
+        return vk::IndexType::eUint8EXT;
     case Maxwell::IndexFormat::UnsignedShort:
         return vk::IndexType::eUint16;
     case Maxwell::IndexFormat::UnsignedInt:

src/video_core/renderer_vulkan/maxwell_to_vk.h

@@ -4,7 +4,6 @@
 
 #pragma once
 
-#include <utility>
 #include "common/common_types.h"
 #include "video_core/engines/maxwell_3d.h"
 #include "video_core/renderer_vulkan/declarations.h"
@@ -23,24 +22,31 @@ vk::Filter Filter(Tegra::Texture::TextureFilter filter);
 
 vk::SamplerMipmapMode MipmapMode(Tegra::Texture::TextureMipmapFilter mipmap_filter);
 
-vk::SamplerAddressMode WrapMode(Tegra::Texture::WrapMode wrap_mode);
+vk::SamplerAddressMode WrapMode(Tegra::Texture::WrapMode wrap_mode,
+                                Tegra::Texture::TextureFilter filter);
 
 vk::CompareOp DepthCompareFunction(Tegra::Texture::DepthCompareFunc depth_compare_func);
 
 } // namespace Sampler
 
-std::pair<vk::Format, bool> SurfaceFormat(const VKDevice& device, FormatType format_type,
-                                          PixelFormat pixel_format);
+struct FormatInfo {
+    vk::Format format;
+    bool attachable;
+    bool storage;
+};
+
+FormatInfo SurfaceFormat(const VKDevice& device, FormatType format_type, PixelFormat pixel_format);
 
 vk::ShaderStageFlagBits ShaderStage(Tegra::Engines::ShaderType stage);
 
-vk::PrimitiveTopology PrimitiveTopology(Maxwell::PrimitiveTopology topology);
+vk::PrimitiveTopology PrimitiveTopology(const VKDevice& device,
+                                        Maxwell::PrimitiveTopology topology);
 
 vk::Format VertexFormat(Maxwell::VertexAttribute::Type type, Maxwell::VertexAttribute::Size size);
 
 vk::CompareOp ComparisonOp(Maxwell::ComparisonOp comparison);
 
-vk::IndexType IndexFormat(Maxwell::IndexFormat index_format);
+vk::IndexType IndexFormat(const VKDevice& device, Maxwell::IndexFormat index_format);
 
 vk::StencilOp StencilOp(Maxwell::StencilOp stencil_op);
 

src/video_core/renderer_vulkan/shaders/blit.frag

@@ -0,0 +1,24 @@
+// Copyright 2019 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+/*
+ * Build instructions:
+ * $ glslangValidator -V $THIS_FILE -o output.spv
+ * $ spirv-opt -O --strip-debug output.spv -o optimized.spv
+ * $ xxd -i optimized.spv
+ *
+ * Then copy that bytecode to the C++ file
+ */
+
+#version 460 core
+
+layout (location = 0) in vec2 frag_tex_coord;
+
+layout (location = 0) out vec4 color;
+
+layout (binding = 1) uniform sampler2D color_texture;
+
+void main() {
+    color = texture(color_texture, frag_tex_coord);
+}

src/video_core/renderer_vulkan/shaders/blit.vert

@@ -0,0 +1,28 @@
+// Copyright 2019 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+/*
+ * Build instructions:
+ * $ glslangValidator -V $THIS_FILE -o output.spv
+ * $ spirv-opt -O --strip-debug output.spv -o optimized.spv
+ * $ xxd -i optimized.spv
+ *
+ * Then copy that bytecode to the C++ file
+ */
+
+#version 460 core
+
+layout (location = 0) in vec2 vert_position;
+layout (location = 1) in vec2 vert_tex_coord;
+
+layout (location = 0) out vec2 frag_tex_coord;
+
+layout (set = 0, binding = 0) uniform MatrixBlock {
+    mat4 modelview_matrix;
+};
+
+void main() {
+    gl_Position = modelview_matrix * vec4(vert_position, 0.0, 1.0);
+    frag_tex_coord = vert_tex_coord;
+}

src/video_core/renderer_vulkan/shaders/quad_array.comp

@@ -0,0 +1,37 @@
+// Copyright 2019 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+/*
+ * Build instructions:
+ * $ glslangValidator -V $THIS_FILE -o output.spv
+ * $ spirv-opt -O --strip-debug output.spv -o optimized.spv
+ * $ xxd -i optimized.spv
+ *
+ * Then copy that bytecode to the C++ file
+ */
+
+#version 460 core
+
+layout (local_size_x = 1024) in;
+
+layout (std430, set = 0, binding = 0) buffer OutputBuffer {
+    uint output_indexes[];
+};
+
+layout (push_constant) uniform PushConstants {
+    uint first;
+};
+
+void main() {
+    uint primitive = gl_GlobalInvocationID.x;
+    if (primitive * 6 >= output_indexes.length()) {
+        return;
+    }
+
+    const uint quad_map[6] = uint[](0, 1, 2, 0, 2, 3);
+    for (uint vertex = 0; vertex < 6; ++vertex) {
+        uint index = first + primitive * 4 + quad_map[vertex];
+        output_indexes[primitive * 6 + vertex] = index;
+    }
+}

src/video_core/renderer_vulkan/shaders/uint8.comp

@@ -0,0 +1,33 @@
+// Copyright 2019 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+/*
+ * Build instructions:
+ * $ glslangValidator -V $THIS_FILE -o output.spv
+ * $ spirv-opt -O --strip-debug output.spv -o optimized.spv
+ * $ xxd -i optimized.spv
+ *
+ * Then copy that bytecode to the C++ file
+ */
+
+#version 460 core
+#extension GL_EXT_shader_16bit_storage : require
+#extension GL_EXT_shader_8bit_storage : require
+
+layout (local_size_x = 1024) in;
+
+layout (std430, set = 0, binding = 0) readonly buffer InputBuffer {
+    uint8_t input_indexes[];
+};
+
+layout (std430, set = 0, binding = 1) writeonly buffer OutputBuffer {
+    uint16_t output_indexes[];
+};
+
+void main() {
+    uint id = gl_GlobalInvocationID.x;
+    if (id < input_indexes.length()) {
+        output_indexes[id] = uint16_t(input_indexes[id]);
+    }
+}

src/video_core/renderer_vulkan/vk_device.cpp

@@ -3,6 +3,7 @@
 // Refer to the license.txt file included.
 
 #include <bitset>
+#include <cstdlib>
 #include <optional>
 #include <set>
 #include <string_view>
@@ -15,6 +16,15 @@ namespace Vulkan {
 
 namespace {
 
+namespace Alternatives {
+
+constexpr std::array Depth24UnormS8Uint = {vk::Format::eD32SfloatS8Uint,
+                                           vk::Format::eD16UnormS8Uint, vk::Format{}};
+constexpr std::array Depth16UnormS8Uint = {vk::Format::eD24UnormS8Uint,
+                                           vk::Format::eD32SfloatS8Uint, vk::Format{}};
+
+} // namespace Alternatives
+
 template <typename T>
 void SetNext(void**& next, T& data) {
     *next = &data;
@@ -22,7 +32,7 @@ void SetNext(void**& next, T& data) {
 }
 
 template <typename T>
-T GetFeatures(vk::PhysicalDevice physical, vk::DispatchLoaderDynamic dldi) {
+T GetFeatures(vk::PhysicalDevice physical, const vk::DispatchLoaderDynamic& dldi) {
     vk::PhysicalDeviceFeatures2 features;
     T extension_features;
     features.pNext = &extension_features;
@@ -30,17 +40,14 @@ T GetFeatures(vk::PhysicalDevice physical, vk::DispatchLoaderDynamic dldi) {
     return extension_features;
 }
 
-} // Anonymous namespace
-
-namespace Alternatives {
-
-constexpr std::array Depth24UnormS8Uint = {vk::Format::eD32SfloatS8Uint,
-                                           vk::Format::eD16UnormS8Uint, vk::Format{}};
-constexpr std::array Depth16UnormS8Uint = {vk::Format::eD24UnormS8Uint,
-                                           vk::Format::eD32SfloatS8Uint, vk::Format{}};
-constexpr std::array Astc = {vk::Format::eA8B8G8R8UnormPack32, vk::Format{}};
-
-} // namespace Alternatives
+template <typename T>
+T GetProperties(vk::PhysicalDevice physical, const vk::DispatchLoaderDynamic& dldi) {
+    vk::PhysicalDeviceProperties2 properties;
+    T extension_properties;
+    properties.pNext = &extension_properties;
+    physical.getProperties2(&properties, dldi);
+    return extension_properties;
+}
 
 constexpr const vk::Format* GetFormatAlternatives(vk::Format format) {
     switch (format) {
@@ -53,8 +60,7 @@ constexpr const vk::Format* GetFormatAlternatives(vk::Format format) {
     }
 }
 
-constexpr vk::FormatFeatureFlags GetFormatFeatures(vk::FormatProperties properties,
-                                                   FormatType format_type) {
+vk::FormatFeatureFlags GetFormatFeatures(vk::FormatProperties properties, FormatType format_type) {
     switch (format_type) {
     case FormatType::Linear:
         return properties.linearTilingFeatures;
@@ -67,11 +73,13 @@ constexpr vk::FormatFeatureFlags GetFormatFeatures(vk::FormatProperties properti
     }
 }
 
+} // Anonymous namespace
+
 VKDevice::VKDevice(const vk::DispatchLoaderDynamic& dldi, vk::PhysicalDevice physical,
                    vk::SurfaceKHR surface)
-    : physical{physical}, format_properties{GetFormatProperties(dldi, physical)} {
+    : physical{physical}, properties{physical.getProperties(dldi)},
+      format_properties{GetFormatProperties(dldi, physical)} {
     SetupFamilies(dldi, surface);
-    SetupProperties(dldi);
     SetupFeatures(dldi);
 }
 
@@ -89,12 +97,22 @@ bool VKDevice::Create(const vk::DispatchLoaderDynamic& dldi, vk::Instance instan
     features.depthClamp = true;
     features.samplerAnisotropy = true;
     features.largePoints = true;
+    features.multiViewport = true;
+    features.depthBiasClamp = true;
+    features.geometryShader = true;
+    features.tessellationShader = true;
+    features.fragmentStoresAndAtomics = true;
+    features.shaderImageGatherExtended = true;
+    features.shaderStorageImageWriteWithoutFormat = true;
     features.textureCompressionASTC_LDR = is_optimal_astc_supported;
 
-    vk::PhysicalDeviceVertexAttributeDivisorFeaturesEXT vertex_divisor;
-    vertex_divisor.vertexAttributeInstanceRateDivisor = true;
-    vertex_divisor.vertexAttributeInstanceRateZeroDivisor = true;
-    SetNext(next, vertex_divisor);
+    vk::PhysicalDevice16BitStorageFeaturesKHR bit16_storage;
+    bit16_storage.uniformAndStorageBuffer16BitAccess = true;
+    SetNext(next, bit16_storage);
+
+    vk::PhysicalDevice8BitStorageFeaturesKHR bit8_storage;
+    bit8_storage.uniformAndStorageBuffer8BitAccess = true;
+    SetNext(next, bit8_storage);
 
     vk::PhysicalDeviceFloat16Int8FeaturesKHR float16_int8;
     if (is_float16_supported) {
@@ -120,6 +138,10 @@ bool VKDevice::Create(const vk::DispatchLoaderDynamic& dldi, vk::Instance instan
         LOG_INFO(Render_Vulkan, "Device doesn't support uint8 indexes");
     }
 
+    if (!ext_depth_range_unrestricted) {
+        LOG_INFO(Render_Vulkan, "Device doesn't support depth range unrestricted");
+    }
+
     vk::DeviceCreateInfo device_ci({}, static_cast<u32>(queue_cis.size()), queue_cis.data(), 0,
                                    nullptr, static_cast<u32>(extensions.size()), extensions.data(),
                                    nullptr);
@@ -135,16 +157,7 @@ bool VKDevice::Create(const vk::DispatchLoaderDynamic& dldi, vk::Instance instan
     logical = UniqueDevice(
         dummy_logical, vk::ObjectDestroy<vk::NoParent, vk::DispatchLoaderDynamic>(nullptr, dld));
 
-    if (khr_driver_properties) {
-        vk::PhysicalDeviceDriverPropertiesKHR driver;
-        vk::PhysicalDeviceProperties2 properties;
-        properties.pNext = &driver;
-        physical.getProperties2(&properties, dld);
-        driver_id = driver.driverID;
-        LOG_INFO(Render_Vulkan, "Driver: {} {}", driver.driverName, driver.driverInfo);
-    } else {
-        LOG_INFO(Render_Vulkan, "Driver: Unknown");
-    }
+    CollectTelemetryParameters();
 
     graphics_queue = logical->getQueue(graphics_family, 0, dld);
     present_queue = logical->getQueue(present_family, 0, dld);
@@ -190,6 +203,18 @@ vk::Format VKDevice::GetSupportedFormat(vk::Format wanted_format,
 
 bool VKDevice::IsOptimalAstcSupported(const vk::PhysicalDeviceFeatures& features,
                                       const vk::DispatchLoaderDynamic& dldi) const {
+    // Disable for now to avoid converting ASTC twice.
+    return false;
+    static constexpr std::array astc_formats = {
+        vk::Format::eAstc4x4SrgbBlock,    vk::Format::eAstc8x8SrgbBlock,
+        vk::Format::eAstc8x5SrgbBlock,    vk::Format::eAstc5x4SrgbBlock,
+        vk::Format::eAstc5x5UnormBlock,   vk::Format::eAstc5x5SrgbBlock,
+        vk::Format::eAstc10x8UnormBlock,  vk::Format::eAstc10x8SrgbBlock,
+        vk::Format::eAstc6x6UnormBlock,   vk::Format::eAstc6x6SrgbBlock,
+        vk::Format::eAstc10x10UnormBlock, vk::Format::eAstc10x10SrgbBlock,
+        vk::Format::eAstc12x12UnormBlock, vk::Format::eAstc12x12SrgbBlock,
+        vk::Format::eAstc8x6UnormBlock,   vk::Format::eAstc8x6SrgbBlock,
+        vk::Format::eAstc6x5UnormBlock,   vk::Format::eAstc6x5SrgbBlock};
     if (!features.textureCompressionASTC_LDR) {
         return false;
     }
@@ -197,12 +222,6 @@ bool VKDevice::IsOptimalAstcSupported(const vk::PhysicalDeviceFeatures& features
         vk::FormatFeatureFlagBits::eSampledImage | vk::FormatFeatureFlagBits::eBlitSrc |
         vk::FormatFeatureFlagBits::eBlitDst | vk::FormatFeatureFlagBits::eTransferSrc |
         vk::FormatFeatureFlagBits::eTransferDst};
-    constexpr std::array astc_formats = {
-        vk::Format::eAstc4x4UnormBlock, vk::Format::eAstc4x4SrgbBlock,
-        vk::Format::eAstc8x8SrgbBlock,  vk::Format::eAstc8x6SrgbBlock,
-        vk::Format::eAstc5x4SrgbBlock,  vk::Format::eAstc5x5UnormBlock,
-        vk::Format::eAstc5x5SrgbBlock,  vk::Format::eAstc10x8UnormBlock,
-        vk::Format::eAstc10x8SrgbBlock};
     for (const auto format : astc_formats) {
         const auto format_properties{physical.getFormatProperties(format, dldi)};
         if (!(format_properties.optimalTilingFeatures & format_feature_usage)) {
@@ -225,11 +244,17 @@ bool VKDevice::IsFormatSupported(vk::Format wanted_format, vk::FormatFeatureFlag
 
 bool VKDevice::IsSuitable(const vk::DispatchLoaderDynamic& dldi, vk::PhysicalDevice physical,
                           vk::SurfaceKHR surface) {
-    LOG_INFO(Render_Vulkan, "{}", physical.getProperties(dldi).deviceName);
     bool is_suitable = true;
 
-    constexpr std::array required_extensions = {VK_KHR_SWAPCHAIN_EXTENSION_NAME,
-                                                VK_EXT_VERTEX_ATTRIBUTE_DIVISOR_EXTENSION_NAME};
+    constexpr std::array required_extensions = {
+        VK_KHR_SWAPCHAIN_EXTENSION_NAME,
+        VK_KHR_16BIT_STORAGE_EXTENSION_NAME,
+        VK_KHR_8BIT_STORAGE_EXTENSION_NAME,
+        VK_KHR_DRIVER_PROPERTIES_EXTENSION_NAME,
+        VK_EXT_VERTEX_ATTRIBUTE_DIVISOR_EXTENSION_NAME,
+        VK_EXT_SHADER_SUBGROUP_BALLOT_EXTENSION_NAME,
+        VK_EXT_SHADER_SUBGROUP_VOTE_EXTENSION_NAME,
+    };
     std::bitset<required_extensions.size()> available_extensions{};
 
     for (const auto& prop : physical.enumerateDeviceExtensionProperties(nullptr, dldi)) {
@@ -246,7 +271,7 @@ bool VKDevice::IsSuitable(const vk::DispatchLoaderDynamic& dldi, vk::PhysicalDev
             if (available_extensions[i]) {
                 continue;
             }
-            LOG_INFO(Render_Vulkan, "Missing required extension: {}", required_extensions[i]);
+            LOG_ERROR(Render_Vulkan, "Missing required extension: {}", required_extensions[i]);
             is_suitable = false;
         }
     }
@@ -263,7 +288,7 @@ bool VKDevice::IsSuitable(const vk::DispatchLoaderDynamic& dldi, vk::PhysicalDev
         has_present |= physical.getSurfaceSupportKHR(i, surface, dldi) != 0;
     }
     if (!has_graphics || !has_present) {
-        LOG_INFO(Render_Vulkan, "Device lacks a graphics and present queue");
+        LOG_ERROR(Render_Vulkan, "Device lacks a graphics and present queue");
         is_suitable = false;
     }
 
@@ -273,8 +298,15 @@ bool VKDevice::IsSuitable(const vk::DispatchLoaderDynamic& dldi, vk::PhysicalDev
 
     constexpr u32 required_ubo_size = 65536;
     if (limits.maxUniformBufferRange < required_ubo_size) {
-        LOG_INFO(Render_Vulkan, "Device UBO size {} is too small, {} is required)",
-                 limits.maxUniformBufferRange, required_ubo_size);
+        LOG_ERROR(Render_Vulkan, "Device UBO size {} is too small, {} is required",
+                  limits.maxUniformBufferRange, required_ubo_size);
+        is_suitable = false;
+    }
+
+    constexpr u32 required_num_viewports = 16;
+    if (limits.maxViewports < required_num_viewports) {
+        LOG_INFO(Render_Vulkan, "Device number of viewports {} is too small, {} is required",
+                 limits.maxViewports, required_num_viewports);
         is_suitable = false;
     }
 
@@ -285,24 +317,32 @@ bool VKDevice::IsSuitable(const vk::DispatchLoaderDynamic& dldi, vk::PhysicalDev
         std::make_pair(features.depthClamp, "depthClamp"),
         std::make_pair(features.samplerAnisotropy, "samplerAnisotropy"),
         std::make_pair(features.largePoints, "largePoints"),
+        std::make_pair(features.multiViewport, "multiViewport"),
+        std::make_pair(features.depthBiasClamp, "depthBiasClamp"),
+        std::make_pair(features.geometryShader, "geometryShader"),
+        std::make_pair(features.tessellationShader, "tessellationShader"),
+        std::make_pair(features.fragmentStoresAndAtomics, "fragmentStoresAndAtomics"),
+        std::make_pair(features.shaderImageGatherExtended, "shaderImageGatherExtended"),
+        std::make_pair(features.shaderStorageImageWriteWithoutFormat,
+                       "shaderStorageImageWriteWithoutFormat"),
     };
     for (const auto& [supported, name] : feature_report) {
         if (supported) {
             continue;
         }
-        LOG_INFO(Render_Vulkan, "Missing required feature: {}", name);
+        LOG_ERROR(Render_Vulkan, "Missing required feature: {}", name);
         is_suitable = false;
     }
 
+    if (!is_suitable) {
+        LOG_ERROR(Render_Vulkan, "{} is not suitable", properties.deviceName);
+    }
+
     return is_suitable;
 }
 
 std::vector<const char*> VKDevice::LoadExtensions(const vk::DispatchLoaderDynamic& dldi) {
     std::vector<const char*> extensions;
-    extensions.reserve(7);
-    extensions.push_back(VK_KHR_SWAPCHAIN_EXTENSION_NAME);
-    extensions.push_back(VK_EXT_VERTEX_ATTRIBUTE_DIVISOR_EXTENSION_NAME);
-
     const auto Test = [&](const vk::ExtensionProperties& extension,
                           std::optional<std::reference_wrapper<bool>> status, const char* name,
                           bool push) {
@@ -317,13 +357,30 @@ std::vector<const char*> VKDevice::LoadExtensions(const vk::DispatchLoaderDynami
         }
     };
 
+    extensions.reserve(13);
+    extensions.push_back(VK_KHR_SWAPCHAIN_EXTENSION_NAME);
+    extensions.push_back(VK_KHR_16BIT_STORAGE_EXTENSION_NAME);
+    extensions.push_back(VK_KHR_8BIT_STORAGE_EXTENSION_NAME);
+    extensions.push_back(VK_KHR_DRIVER_PROPERTIES_EXTENSION_NAME);
+    extensions.push_back(VK_EXT_VERTEX_ATTRIBUTE_DIVISOR_EXTENSION_NAME);
+    extensions.push_back(VK_EXT_SHADER_SUBGROUP_BALLOT_EXTENSION_NAME);
+    extensions.push_back(VK_EXT_SHADER_SUBGROUP_VOTE_EXTENSION_NAME);
+
+    [[maybe_unused]] const bool nsight =
+        std::getenv("NVTX_INJECTION64_PATH") || std::getenv("NSIGHT_LAUNCHED");
     bool khr_shader_float16_int8{};
+    bool ext_subgroup_size_control{};
     for (const auto& extension : physical.enumerateDeviceExtensionProperties(nullptr, dldi)) {
         Test(extension, khr_uniform_buffer_standard_layout,
              VK_KHR_UNIFORM_BUFFER_STANDARD_LAYOUT_EXTENSION_NAME, true);
-        Test(extension, ext_index_type_uint8, VK_EXT_INDEX_TYPE_UINT8_EXTENSION_NAME, true);
-        Test(extension, khr_driver_properties, VK_KHR_DRIVER_PROPERTIES_EXTENSION_NAME, true);
         Test(extension, khr_shader_float16_int8, VK_KHR_SHADER_FLOAT16_INT8_EXTENSION_NAME, false);
+        Test(extension, ext_depth_range_unrestricted,
+             VK_EXT_DEPTH_RANGE_UNRESTRICTED_EXTENSION_NAME, true);
+        Test(extension, ext_index_type_uint8, VK_EXT_INDEX_TYPE_UINT8_EXTENSION_NAME, true);
+        Test(extension, ext_shader_viewport_index_layer,
+             VK_EXT_SHADER_VIEWPORT_INDEX_LAYER_EXTENSION_NAME, true);
+        Test(extension, ext_subgroup_size_control, VK_EXT_SUBGROUP_SIZE_CONTROL_EXTENSION_NAME,
+             false);
     }
 
     if (khr_shader_float16_int8) {
@@ -332,6 +389,23 @@ std::vector<const char*> VKDevice::LoadExtensions(const vk::DispatchLoaderDynami
         extensions.push_back(VK_KHR_SHADER_FLOAT16_INT8_EXTENSION_NAME);
     }
 
+    if (ext_subgroup_size_control) {
+        const auto features =
+            GetFeatures<vk::PhysicalDeviceSubgroupSizeControlFeaturesEXT>(physical, dldi);
+        const auto properties =
+            GetProperties<vk::PhysicalDeviceSubgroupSizeControlPropertiesEXT>(physical, dldi);
+
+        is_warp_potentially_bigger = properties.maxSubgroupSize > GuestWarpSize;
+
+        if (features.subgroupSizeControl && properties.minSubgroupSize <= GuestWarpSize &&
+            properties.maxSubgroupSize >= GuestWarpSize) {
+            extensions.push_back(VK_EXT_SUBGROUP_SIZE_CONTROL_EXTENSION_NAME);
+            guest_warp_stages = properties.requiredSubgroupSizeStages;
+        }
+    } else {
+        is_warp_potentially_bigger = true;
+    }
+
     return extensions;
 }
 
@@ -358,19 +432,23 @@ void VKDevice::SetupFamilies(const vk::DispatchLoaderDynamic& dldi, vk::SurfaceK
     present_family = *present_family_;
 }
 
-void VKDevice::SetupProperties(const vk::DispatchLoaderDynamic& dldi) {
-    const auto props = physical.getProperties(dldi);
-    device_type = props.deviceType;
-    uniform_buffer_alignment = static_cast<u64>(props.limits.minUniformBufferOffsetAlignment);
-    storage_buffer_alignment = static_cast<u64>(props.limits.minStorageBufferOffsetAlignment);
-    max_storage_buffer_range = static_cast<u64>(props.limits.maxStorageBufferRange);
-}
-
 void VKDevice::SetupFeatures(const vk::DispatchLoaderDynamic& dldi) {
     const auto supported_features{physical.getFeatures(dldi)};
     is_optimal_astc_supported = IsOptimalAstcSupported(supported_features, dldi);
 }
 
+void VKDevice::CollectTelemetryParameters() {
+    const auto driver = GetProperties<vk::PhysicalDeviceDriverPropertiesKHR>(physical, dld);
+    driver_id = driver.driverID;
+    vendor_name = driver.driverName;
+
+    const auto extensions = physical.enumerateDeviceExtensionProperties(nullptr, dld);
+    reported_extensions.reserve(std::size(extensions));
+    for (const auto& extension : extensions) {
+        reported_extensions.push_back(extension.extensionName);
+    }
+}
+
 std::vector<vk::DeviceQueueCreateInfo> VKDevice::GetDeviceQueueCreateInfos() const {
     static const float QUEUE_PRIORITY = 1.0f;
 
@@ -385,50 +463,70 @@ std::vector<vk::DeviceQueueCreateInfo> VKDevice::GetDeviceQueueCreateInfos() con
 
 std::unordered_map<vk::Format, vk::FormatProperties> VKDevice::GetFormatProperties(
     const vk::DispatchLoaderDynamic& dldi, vk::PhysicalDevice physical) {
-    constexpr std::array formats{vk::Format::eA8B8G8R8UnormPack32,
-                                 vk::Format::eA8B8G8R8SnormPack32,
-                                 vk::Format::eA8B8G8R8SrgbPack32,
-                                 vk::Format::eB5G6R5UnormPack16,
-                                 vk::Format::eA2B10G10R10UnormPack32,
-                                 vk::Format::eR32G32B32A32Sfloat,
-                                 vk::Format::eR16G16B16A16Uint,
-                                 vk::Format::eR16G16Unorm,
-                                 vk::Format::eR16G16Snorm,
-                                 vk::Format::eR16G16Sfloat,
-                                 vk::Format::eR16Unorm,
-                                 vk::Format::eR8G8B8A8Srgb,
-                                 vk::Format::eR8G8Unorm,
-                                 vk::Format::eR8G8Snorm,
-                                 vk::Format::eR8Unorm,
-                                 vk::Format::eB10G11R11UfloatPack32,
-                                 vk::Format::eR32Sfloat,
-                                 vk::Format::eR16Sfloat,
-                                 vk::Format::eR16G16B16A16Sfloat,
-                                 vk::Format::eB8G8R8A8Unorm,
-                                 vk::Format::eD32Sfloat,
-                                 vk::Format::eD16Unorm,
-                                 vk::Format::eD16UnormS8Uint,
-                                 vk::Format::eD24UnormS8Uint,
-                                 vk::Format::eD32SfloatS8Uint,
-                                 vk::Format::eBc1RgbaUnormBlock,
-                                 vk::Format::eBc2UnormBlock,
-                                 vk::Format::eBc3UnormBlock,
-                                 vk::Format::eBc4UnormBlock,
-                                 vk::Format::eBc5UnormBlock,
-                                 vk::Format::eBc5SnormBlock,
-                                 vk::Format::eBc7UnormBlock,
-                                 vk::Format::eBc1RgbaSrgbBlock,
-                                 vk::Format::eBc3SrgbBlock,
-                                 vk::Format::eBc7SrgbBlock,
-                                 vk::Format::eAstc4x4UnormBlock,
-                                 vk::Format::eAstc4x4SrgbBlock,
-                                 vk::Format::eAstc8x8SrgbBlock,
-                                 vk::Format::eAstc8x6SrgbBlock,
-                                 vk::Format::eAstc5x4SrgbBlock,
-                                 vk::Format::eAstc5x5UnormBlock,
-                                 vk::Format::eAstc5x5SrgbBlock,
-                                 vk::Format::eAstc10x8UnormBlock,
-                                 vk::Format::eAstc10x8SrgbBlock};
+    static constexpr std::array formats{vk::Format::eA8B8G8R8UnormPack32,
+                                        vk::Format::eA8B8G8R8SnormPack32,
+                                        vk::Format::eA8B8G8R8SrgbPack32,
+                                        vk::Format::eB5G6R5UnormPack16,
+                                        vk::Format::eA2B10G10R10UnormPack32,
+                                        vk::Format::eA1R5G5B5UnormPack16,
+                                        vk::Format::eR32G32B32A32Sfloat,
+                                        vk::Format::eR32G32B32A32Uint,
+                                        vk::Format::eR32G32Sfloat,
+                                        vk::Format::eR32G32Uint,
+                                        vk::Format::eR16G16B16A16Uint,
+                                        vk::Format::eR16G16B16A16Unorm,
+                                        vk::Format::eR16G16Unorm,
+                                        vk::Format::eR16G16Snorm,
+                                        vk::Format::eR16G16Sfloat,
+                                        vk::Format::eR16Unorm,
+                                        vk::Format::eR8G8B8A8Srgb,
+                                        vk::Format::eR8G8Unorm,
+                                        vk::Format::eR8G8Snorm,
+                                        vk::Format::eR8Unorm,
+                                        vk::Format::eR8Uint,
+                                        vk::Format::eB10G11R11UfloatPack32,
+                                        vk::Format::eR32Sfloat,
+                                        vk::Format::eR32Uint,
+                                        vk::Format::eR16Sfloat,
+                                        vk::Format::eR16G16B16A16Sfloat,
+                                        vk::Format::eB8G8R8A8Unorm,
+                                        vk::Format::eR4G4B4A4UnormPack16,
+                                        vk::Format::eD32Sfloat,
+                                        vk::Format::eD16Unorm,
+                                        vk::Format::eD16UnormS8Uint,
+                                        vk::Format::eD24UnormS8Uint,
+                                        vk::Format::eD32SfloatS8Uint,
+                                        vk::Format::eBc1RgbaUnormBlock,
+                                        vk::Format::eBc2UnormBlock,
+                                        vk::Format::eBc3UnormBlock,
+                                        vk::Format::eBc4UnormBlock,
+                                        vk::Format::eBc5UnormBlock,
+                                        vk::Format::eBc5SnormBlock,
+                                        vk::Format::eBc7UnormBlock,
+                                        vk::Format::eBc6HUfloatBlock,
+                                        vk::Format::eBc6HSfloatBlock,
+                                        vk::Format::eBc1RgbaSrgbBlock,
+                                        vk::Format::eBc3SrgbBlock,
+                                        vk::Format::eBc7SrgbBlock,
+                                        vk::Format::eAstc4x4SrgbBlock,
+                                        vk::Format::eAstc8x8SrgbBlock,
+                                        vk::Format::eAstc8x5SrgbBlock,
+                                        vk::Format::eAstc5x4SrgbBlock,
+                                        vk::Format::eAstc5x5UnormBlock,
+                                        vk::Format::eAstc5x5SrgbBlock,
+                                        vk::Format::eAstc10x8UnormBlock,
+                                        vk::Format::eAstc10x8SrgbBlock,
+                                        vk::Format::eAstc6x6UnormBlock,
+                                        vk::Format::eAstc6x6SrgbBlock,
+                                        vk::Format::eAstc10x10UnormBlock,
+                                        vk::Format::eAstc10x10SrgbBlock,
+                                        vk::Format::eAstc12x12UnormBlock,
+                                        vk::Format::eAstc12x12SrgbBlock,
+                                        vk::Format::eAstc8x6UnormBlock,
+                                        vk::Format::eAstc8x6SrgbBlock,
+                                        vk::Format::eAstc6x5UnormBlock,
+                                        vk::Format::eAstc6x5SrgbBlock,
+                                        vk::Format::eE5B9G9R9UfloatPack32};
     std::unordered_map<vk::Format, vk::FormatProperties> format_properties;
     for (const auto format : formats) {
         format_properties.emplace(format, physical.getFormatProperties(format, dldi));

src/video_core/renderer_vulkan/vk_device.h

@@ -4,6 +4,8 @@
 
 #pragma once
 
+#include <string>
+#include <string_view>
 #include <unordered_map>
 #include <vector>
 #include "common/common_types.h"
@@ -14,6 +16,9 @@ namespace Vulkan {
 /// Format usage descriptor.
 enum class FormatType { Linear, Optimal, Buffer };
 
+/// Subgroup size of the guest emulated hardware (Nvidia has 32 threads per subgroup).
+const u32 GuestWarpSize = 32;
+
 /// Handles data specific to a physical device.
 class VKDevice final {
 public:
@@ -71,7 +76,22 @@ public:
 
     /// Returns true if the device is integrated with the host CPU.
     bool IsIntegrated() const {
-        return device_type == vk::PhysicalDeviceType::eIntegratedGpu;
+        return properties.deviceType == vk::PhysicalDeviceType::eIntegratedGpu;
+    }
+
+    /// Returns the current Vulkan API version provided in Vulkan-formatted version numbers.
+    u32 GetApiVersion() const {
+        return properties.apiVersion;
+    }
+
+    /// Returns the current driver version provided in Vulkan-formatted version numbers.
+    u32 GetDriverVersion() const {
+        return properties.driverVersion;
+    }
+
+    /// Returns the device name.
+    std::string_view GetModelName() const {
+        return properties.deviceName;
     }
 
     /// Returns the driver ID.
@@ -80,18 +100,23 @@ public:
     }
 
     /// Returns uniform buffer alignment requeriment.
-    u64 GetUniformBufferAlignment() const {
-        return uniform_buffer_alignment;
+    vk::DeviceSize GetUniformBufferAlignment() const {
+        return properties.limits.minUniformBufferOffsetAlignment;
     }
 
     /// Returns storage alignment requeriment.
-    u64 GetStorageBufferAlignment() const {
-        return storage_buffer_alignment;
+    vk::DeviceSize GetStorageBufferAlignment() const {
+        return properties.limits.minStorageBufferOffsetAlignment;
     }
 
     /// Returns the maximum range for storage buffers.
-    u64 GetMaxStorageBufferRange() const {
-        return max_storage_buffer_range;
+    vk::DeviceSize GetMaxStorageBufferRange() const {
+        return properties.limits.maxStorageBufferRange;
+    }
+
+    /// Returns the maximum size for push constants.
+    vk::DeviceSize GetMaxPushConstantsSize() const {
+        return properties.limits.maxPushConstantsSize;
     }
 
     /// Returns true if ASTC is natively supported.
@@ -104,6 +129,16 @@ public:
         return is_float16_supported;
     }
 
+    /// Returns true if the device warp size can potentially be bigger than guest's warp size.
+    bool IsWarpSizePotentiallyBiggerThanGuest() const {
+        return is_warp_potentially_bigger;
+    }
+
+    /// Returns true if the device can be forced to use the guest warp size.
+    bool IsGuestWarpSizeSupported(vk::ShaderStageFlagBits stage) const {
+        return (guest_warp_stages & stage) != vk::ShaderStageFlags{};
+    }
+
     /// Returns true if the device supports VK_EXT_scalar_block_layout.
     bool IsKhrUniformBufferStandardLayoutSupported() const {
         return khr_uniform_buffer_standard_layout;
@@ -114,6 +149,26 @@ public:
         return ext_index_type_uint8;
     }
 
+    /// Returns true if the device supports VK_EXT_depth_range_unrestricted.
+    bool IsExtDepthRangeUnrestrictedSupported() const {
+        return ext_depth_range_unrestricted;
+    }
+
+    /// Returns true if the device supports VK_EXT_shader_viewport_index_layer.
+    bool IsExtShaderViewportIndexLayerSupported() const {
+        return ext_shader_viewport_index_layer;
+    }
+
+    /// Returns the vendor name reported from Vulkan.
+    std::string_view GetVendorName() const {
+        return vendor_name;
+    }
+
+    /// Returns the list of available extensions.
+    const std::vector<std::string>& GetAvailableExtensions() const {
+        return reported_extensions;
+    }
+
     /// Checks if the physical device is suitable.
     static bool IsSuitable(const vk::DispatchLoaderDynamic& dldi, vk::PhysicalDevice physical,
                            vk::SurfaceKHR surface);
@@ -125,12 +180,12 @@ private:
     /// Sets up queue families.
     void SetupFamilies(const vk::DispatchLoaderDynamic& dldi, vk::SurfaceKHR surface);
 
-    /// Sets up device properties.
-    void SetupProperties(const vk::DispatchLoaderDynamic& dldi);
-
     /// Sets up device features.
     void SetupFeatures(const vk::DispatchLoaderDynamic& dldi);
 
+    /// Collects telemetry information from the device.
+    void CollectTelemetryParameters();
+
     /// Returns a list of queue initialization descriptors.
     std::vector<vk::DeviceQueueCreateInfo> GetDeviceQueueCreateInfos() const;
 
@@ -148,23 +203,28 @@ private:
 
     const vk::PhysicalDevice physical;         ///< Physical device.
     vk::DispatchLoaderDynamic dld;             ///< Device function pointers.
+    vk::PhysicalDeviceProperties properties;   ///< Device properties.
     UniqueDevice logical;                      ///< Logical device.
     vk::Queue graphics_queue;                  ///< Main graphics queue.
     vk::Queue present_queue;                   ///< Main present queue.
     u32 graphics_family{};                     ///< Main graphics queue family index.
     u32 present_family{};                      ///< Main present queue family index.
-    vk::PhysicalDeviceType device_type;        ///< Physical device type.
     vk::DriverIdKHR driver_id{};               ///< Driver ID.
-    u64 uniform_buffer_alignment{};            ///< Uniform buffer alignment requeriment.
-    u64 storage_buffer_alignment{};            ///< Storage buffer alignment requeriment.
-    u64 max_storage_buffer_range{};            ///< Max storage buffer size.
+    vk::ShaderStageFlags guest_warp_stages{};  ///< Stages where the guest warp size can be forced.
     bool is_optimal_astc_supported{};          ///< Support for native ASTC.
     bool is_float16_supported{};               ///< Support for float16 arithmetics.
+    bool is_warp_potentially_bigger{};         ///< Host warp size can be bigger than guest.
     bool khr_uniform_buffer_standard_layout{}; ///< Support for std430 on UBOs.
     bool ext_index_type_uint8{};               ///< Support for VK_EXT_index_type_uint8.
-    bool khr_driver_properties{};              ///< Support for VK_KHR_driver_properties.
-    std::unordered_map<vk::Format, vk::FormatProperties>
-        format_properties; ///< Format properties dictionary.
+    bool ext_depth_range_unrestricted{};       ///< Support for VK_EXT_depth_range_unrestricted.
+    bool ext_shader_viewport_index_layer{};    ///< Support for VK_EXT_shader_viewport_index_layer.
+
+    // Telemetry parameters
+    std::string vendor_name;                      ///< Device's driver name.
+    std::vector<std::string> reported_extensions; ///< Reported Vulkan extensions.
+
+    /// Format properties dictionary.
+    std::unordered_map<vk::Format, vk::FormatProperties> format_properties;
 };
 
 } // namespace Vulkan

src/video_core/renderer_vulkan/vk_sampler_cache.cpp

@@ -46,9 +46,10 @@ UniqueSampler VKSamplerCache::CreateSampler(const Tegra::Texture::TSCEntry& tsc)
         {}, MaxwellToVK::Sampler::Filter(tsc.mag_filter),
         MaxwellToVK::Sampler::Filter(tsc.min_filter),
         MaxwellToVK::Sampler::MipmapMode(tsc.mipmap_filter),
-        MaxwellToVK::Sampler::WrapMode(tsc.wrap_u), MaxwellToVK::Sampler::WrapMode(tsc.wrap_v),
-        MaxwellToVK::Sampler::WrapMode(tsc.wrap_p), tsc.GetLodBias(), has_anisotropy,
-        max_anisotropy, tsc.depth_compare_enabled,
+        MaxwellToVK::Sampler::WrapMode(tsc.wrap_u, tsc.mag_filter),
+        MaxwellToVK::Sampler::WrapMode(tsc.wrap_v, tsc.mag_filter),
+        MaxwellToVK::Sampler::WrapMode(tsc.wrap_p, tsc.mag_filter), tsc.GetLodBias(),
+        has_anisotropy, max_anisotropy, tsc.depth_compare_enabled,
         MaxwellToVK::Sampler::DepthCompareFunction(tsc.depth_compare_func), tsc.GetMinLod(),
         tsc.GetMaxLod(), vk_border_color.value_or(vk::BorderColor::eFloatTransparentBlack),
         unnormalized_coords);

src/video_core/renderer_vulkan/vk_scheduler.cpp

@@ -3,7 +3,7 @@
 // Refer to the license.txt file included.
 
 #include "common/assert.h"
-#include "common/logging/log.h"
+#include "common/microprofile.h"
 #include "video_core/renderer_vulkan/declarations.h"
 #include "video_core/renderer_vulkan/vk_device.h"
 #include "video_core/renderer_vulkan/vk_resource_manager.h"
@@ -11,46 +11,172 @@
 
 namespace Vulkan {
 
-VKScheduler::VKScheduler(const VKDevice& device, VKResourceManager& resource_manager)
-    : device{device}, resource_manager{resource_manager} {
-    next_fence = &resource_manager.CommitFence();
-    AllocateNewContext();
+MICROPROFILE_DECLARE(Vulkan_WaitForWorker);
+
+void VKScheduler::CommandChunk::ExecuteAll(vk::CommandBuffer cmdbuf,
+                                           const vk::DispatchLoaderDynamic& dld) {
+    auto command = first;
+    while (command != nullptr) {
+        auto next = command->GetNext();
+        command->Execute(cmdbuf, dld);
+        command->~Command();
+        command = next;
+    }
+
+    command_offset = 0;
+    first = nullptr;
+    last = nullptr;
 }
 
-VKScheduler::~VKScheduler() = default;
+VKScheduler::VKScheduler(const VKDevice& device, VKResourceManager& resource_manager)
+    : device{device}, resource_manager{resource_manager}, next_fence{
+                                                              &resource_manager.CommitFence()} {
+    AcquireNewChunk();
+    AllocateNewContext();
+    worker_thread = std::thread(&VKScheduler::WorkerThread, this);
+}
+
+VKScheduler::~VKScheduler() {
+    quit = true;
+    cv.notify_all();
+    worker_thread.join();
+}
 
 void VKScheduler::Flush(bool release_fence, vk::Semaphore semaphore) {
     SubmitExecution(semaphore);
-    if (release_fence)
+    if (release_fence) {
         current_fence->Release();
+    }
     AllocateNewContext();
 }
 
 void VKScheduler::Finish(bool release_fence, vk::Semaphore semaphore) {
     SubmitExecution(semaphore);
     current_fence->Wait();
-    if (release_fence)
+    if (release_fence) {
         current_fence->Release();
+    }
     AllocateNewContext();
 }
 
+void VKScheduler::WaitWorker() {
+    MICROPROFILE_SCOPE(Vulkan_WaitForWorker);
+    DispatchWork();
+
+    bool finished = false;
+    do {
+        cv.notify_all();
+        std::unique_lock lock{mutex};
+        finished = chunk_queue.Empty();
+    } while (!finished);
+}
+
+void VKScheduler::DispatchWork() {
+    if (chunk->Empty()) {
+        return;
+    }
+    chunk_queue.Push(std::move(chunk));
+    cv.notify_all();
+    AcquireNewChunk();
+}
+
+void VKScheduler::RequestRenderpass(const vk::RenderPassBeginInfo& renderpass_bi) {
+    if (state.renderpass && renderpass_bi == *state.renderpass) {
+        return;
+    }
+    const bool end_renderpass = state.renderpass.has_value();
+    state.renderpass = renderpass_bi;
+    Record([renderpass_bi, end_renderpass](auto cmdbuf, auto& dld) {
+        if (end_renderpass) {
+            cmdbuf.endRenderPass(dld);
+        }
+        cmdbuf.beginRenderPass(renderpass_bi, vk::SubpassContents::eInline, dld);
+    });
+}
+
+void VKScheduler::RequestOutsideRenderPassOperationContext() {
+    EndRenderPass();
+}
+
+void VKScheduler::BindGraphicsPipeline(vk::Pipeline pipeline) {
+    if (state.graphics_pipeline == pipeline) {
+        return;
+    }
+    state.graphics_pipeline = pipeline;
+    Record([pipeline](auto cmdbuf, auto& dld) {
+        cmdbuf.bindPipeline(vk::PipelineBindPoint::eGraphics, pipeline, dld);
+    });
+}
+
+void VKScheduler::WorkerThread() {
+    std::unique_lock lock{mutex};
+    do {
+        cv.wait(lock, [this] { return !chunk_queue.Empty() || quit; });
+        if (quit) {
+            continue;
+        }
+        auto extracted_chunk = std::move(chunk_queue.Front());
+        chunk_queue.Pop();
+        extracted_chunk->ExecuteAll(current_cmdbuf, device.GetDispatchLoader());
+        chunk_reserve.Push(std::move(extracted_chunk));
+    } while (!quit);
+}
+
 void VKScheduler::SubmitExecution(vk::Semaphore semaphore) {
+    EndPendingOperations();
+    InvalidateState();
+    WaitWorker();
+
+    std::unique_lock lock{mutex};
+
+    const auto queue = device.GetGraphicsQueue();
     const auto& dld = device.GetDispatchLoader();
     current_cmdbuf.end(dld);
 
-    const auto queue = device.GetGraphicsQueue();
-    const vk::SubmitInfo submit_info(0, nullptr, nullptr, 1, &current_cmdbuf, semaphore ? 1u : 0u,
+    const vk::SubmitInfo submit_info(0, nullptr, nullptr, 1, &current_cmdbuf, semaphore ? 1U : 0U,
                                      &semaphore);
-    queue.submit({submit_info}, *current_fence, dld);
+    queue.submit({submit_info}, static_cast<vk::Fence>(*current_fence), dld);
 }
 
 void VKScheduler::AllocateNewContext() {
+    std::unique_lock lock{mutex};
     current_fence = next_fence;
-    current_cmdbuf = resource_manager.CommitCommandBuffer(*current_fence);
     next_fence = &resource_manager.CommitFence();
 
-    const auto& dld = device.GetDispatchLoader();
-    current_cmdbuf.begin({vk::CommandBufferUsageFlagBits::eOneTimeSubmit}, dld);
+    current_cmdbuf = resource_manager.CommitCommandBuffer(*current_fence);
+    current_cmdbuf.begin({vk::CommandBufferUsageFlagBits::eOneTimeSubmit},
+                         device.GetDispatchLoader());
+}
+
+void VKScheduler::InvalidateState() {
+    state.graphics_pipeline = nullptr;
+    state.viewports = false;
+    state.scissors = false;
+    state.depth_bias = false;
+    state.blend_constants = false;
+    state.depth_bounds = false;
+    state.stencil_values = false;
+}
+
+void VKScheduler::EndPendingOperations() {
+    EndRenderPass();
+}
+
+void VKScheduler::EndRenderPass() {
+    if (!state.renderpass) {
+        return;
+    }
+    state.renderpass = std::nullopt;
+    Record([](auto cmdbuf, auto& dld) { cmdbuf.endRenderPass(dld); });
+}
+
+void VKScheduler::AcquireNewChunk() {
+    if (chunk_reserve.Empty()) {
+        chunk = std::make_unique<CommandChunk>();
+        return;
+    }
+    chunk = std::move(chunk_reserve.Front());
+    chunk_reserve.Pop();
 }
 
 } // namespace Vulkan

src/video_core/renderer_vulkan/vk_scheduler.h

@@ -4,7 +4,14 @@
 
 #pragma once
 
+#include <condition_variable>
+#include <memory>
+#include <optional>
+#include <stack>
+#include <thread>
+#include <utility>
 #include "common/common_types.h"
+#include "common/threadsafe_queue.h"
 #include "video_core/renderer_vulkan/declarations.h"
 
 namespace Vulkan {
@@ -30,23 +37,6 @@ private:
     VKFence* const& fence;
 };
 
-class VKCommandBufferView {
-public:
-    VKCommandBufferView() = default;
-    VKCommandBufferView(const vk::CommandBuffer& cmdbuf) : cmdbuf{cmdbuf} {}
-
-    const vk::CommandBuffer* operator->() const noexcept {
-        return &cmdbuf;
-    }
-
-    operator vk::CommandBuffer() const noexcept {
-        return cmdbuf;
-    }
-
-private:
-    const vk::CommandBuffer& cmdbuf;
-};
-
 /// The scheduler abstracts command buffer and fence management with an interface that's able to do
 /// OpenGL-like operations on Vulkan command buffers.
 class VKScheduler {
@@ -54,32 +44,190 @@ public:
     explicit VKScheduler(const VKDevice& device, VKResourceManager& resource_manager);
     ~VKScheduler();
 
-    /// Gets a reference to the current fence.
-    VKFenceView GetFence() const {
-        return current_fence;
-    }
-
-    /// Gets a reference to the current command buffer.
-    VKCommandBufferView GetCommandBuffer() const {
-        return current_cmdbuf;
-    }
-
     /// Sends the current execution context to the GPU.
     void Flush(bool release_fence = true, vk::Semaphore semaphore = nullptr);
 
     /// Sends the current execution context to the GPU and waits for it to complete.
     void Finish(bool release_fence = true, vk::Semaphore semaphore = nullptr);
 
+    /// Waits for the worker thread to finish executing everything. After this function returns it's
+    /// safe to touch worker resources.
+    void WaitWorker();
+
+    /// Sends currently recorded work to the worker thread.
+    void DispatchWork();
+
+    /// Requests to begin a renderpass.
+    void RequestRenderpass(const vk::RenderPassBeginInfo& renderpass_bi);
+
+    /// Requests the current executino context to be able to execute operations only allowed outside
+    /// of a renderpass.
+    void RequestOutsideRenderPassOperationContext();
+
+    /// Binds a pipeline to the current execution context.
+    void BindGraphicsPipeline(vk::Pipeline pipeline);
+
+    /// Returns true when viewports have been set in the current command buffer.
+    bool TouchViewports() {
+        return std::exchange(state.viewports, true);
+    }
+
+    /// Returns true when scissors have been set in the current command buffer.
+    bool TouchScissors() {
+        return std::exchange(state.scissors, true);
+    }
+
+    /// Returns true when depth bias have been set in the current command buffer.
+    bool TouchDepthBias() {
+        return std::exchange(state.depth_bias, true);
+    }
+
+    /// Returns true when blend constants have been set in the current command buffer.
+    bool TouchBlendConstants() {
+        return std::exchange(state.blend_constants, true);
+    }
+
+    /// Returns true when depth bounds have been set in the current command buffer.
+    bool TouchDepthBounds() {
+        return std::exchange(state.depth_bounds, true);
+    }
+
+    /// Returns true when stencil values have been set in the current command buffer.
+    bool TouchStencilValues() {
+        return std::exchange(state.stencil_values, true);
+    }
+
+    /// Send work to a separate thread.
+    template <typename T>
+    void Record(T&& command) {
+        if (chunk->Record(command)) {
+            return;
+        }
+        DispatchWork();
+        (void)chunk->Record(command);
+    }
+
+    /// Gets a reference to the current fence.
+    VKFenceView GetFence() const {
+        return current_fence;
+    }
+
 private:
+    class Command {
+    public:
+        virtual ~Command() = default;
+
+        virtual void Execute(vk::CommandBuffer cmdbuf,
+                             const vk::DispatchLoaderDynamic& dld) const = 0;
+
+        Command* GetNext() const {
+            return next;
+        }
+
+        void SetNext(Command* next_) {
+            next = next_;
+        }
+
+    private:
+        Command* next = nullptr;
+    };
+
+    template <typename T>
+    class TypedCommand final : public Command {
+    public:
+        explicit TypedCommand(T&& command) : command{std::move(command)} {}
+        ~TypedCommand() override = default;
+
+        TypedCommand(TypedCommand&&) = delete;
+        TypedCommand& operator=(TypedCommand&&) = delete;
+
+        void Execute(vk::CommandBuffer cmdbuf,
+                     const vk::DispatchLoaderDynamic& dld) const override {
+            command(cmdbuf, dld);
+        }
+
+    private:
+        T command;
+    };
+
+    class CommandChunk final {
+    public:
+        void ExecuteAll(vk::CommandBuffer cmdbuf, const vk::DispatchLoaderDynamic& dld);
+
+        template <typename T>
+        bool Record(T& command) {
+            using FuncType = TypedCommand<T>;
+            static_assert(sizeof(FuncType) < sizeof(data), "Lambda is too large");
+
+            if (command_offset > sizeof(data) - sizeof(FuncType)) {
+                return false;
+            }
+
+            Command* current_last = last;
+
+            last = new (data.data() + command_offset) FuncType(std::move(command));
+
+            if (current_last) {
+                current_last->SetNext(last);
+            } else {
+                first = last;
+            }
+
+            command_offset += sizeof(FuncType);
+            return true;
+        }
+
+        bool Empty() const {
+            return command_offset == 0;
+        }
+
+    private:
+        Command* first = nullptr;
+        Command* last = nullptr;
+
+        std::size_t command_offset = 0;
+        std::array<u8, 0x8000> data{};
+    };
+
+    void WorkerThread();
+
     void SubmitExecution(vk::Semaphore semaphore);
 
     void AllocateNewContext();
 
+    void InvalidateState();
+
+    void EndPendingOperations();
+
+    void EndRenderPass();
+
+    void AcquireNewChunk();
+
     const VKDevice& device;
     VKResourceManager& resource_manager;
     vk::CommandBuffer current_cmdbuf;
     VKFence* current_fence = nullptr;
     VKFence* next_fence = nullptr;
+
+    struct State {
+        std::optional<vk::RenderPassBeginInfo> renderpass;
+        vk::Pipeline graphics_pipeline;
+        bool viewports = false;
+        bool scissors = false;
+        bool depth_bias = false;
+        bool blend_constants = false;
+        bool depth_bounds = false;
+        bool stencil_values = false;
+    } state;
+
+    std::unique_ptr<CommandChunk> chunk;
+    std::thread worker_thread;
+
+    Common::SPSCQueue<std::unique_ptr<CommandChunk>> chunk_queue;
+    Common::SPSCQueue<std::unique_ptr<CommandChunk>> chunk_reserve;
+    std::mutex mutex;
+    std::condition_variable cv;
+    bool quit = false;
 };
 
 } // namespace Vulkan

src/video_core/renderer_vulkan/vk_shader_decompiler.h

@@ -5,29 +5,28 @@
 #pragma once
 
 #include <array>
+#include <bitset>
 #include <memory>
 #include <set>
+#include <type_traits>
 #include <utility>
 #include <vector>
 
-#include <sirit/sirit.h>
-
 #include "common/common_types.h"
 #include "video_core/engines/maxwell_3d.h"
+#include "video_core/engines/shader_type.h"
 #include "video_core/shader/shader_ir.h"
 
-namespace VideoCommon::Shader {
-class ShaderIR;
-}
-
 namespace Vulkan {
 class VKDevice;
 }
 
-namespace Vulkan::VKShader {
+namespace Vulkan {
 
 using Maxwell = Tegra::Engines::Maxwell3D::Regs;
+using TexelBufferEntry = VideoCommon::Shader::Sampler;
 using SamplerEntry = VideoCommon::Shader::Sampler;
+using ImageEntry = VideoCommon::Shader::Image;
 
 constexpr u32 DESCRIPTOR_SET = 0;
 
@@ -46,39 +45,74 @@ private:
 
 class GlobalBufferEntry {
 public:
-    explicit GlobalBufferEntry(u32 cbuf_index, u32 cbuf_offset)
-        : cbuf_index{cbuf_index}, cbuf_offset{cbuf_offset} {}
+    constexpr explicit GlobalBufferEntry(u32 cbuf_index, u32 cbuf_offset, bool is_written)
+        : cbuf_index{cbuf_index}, cbuf_offset{cbuf_offset}, is_written{is_written} {}
 
-    u32 GetCbufIndex() const {
+    constexpr u32 GetCbufIndex() const {
         return cbuf_index;
     }
 
-    u32 GetCbufOffset() const {
+    constexpr u32 GetCbufOffset() const {
         return cbuf_offset;
     }
 
+    constexpr bool IsWritten() const {
+        return is_written;
+    }
+
 private:
     u32 cbuf_index{};
     u32 cbuf_offset{};
+    bool is_written{};
 };
 
 struct ShaderEntries {
-    u32 const_buffers_base_binding{};
-    u32 global_buffers_base_binding{};
-    u32 samplers_base_binding{};
+    u32 NumBindings() const {
+        return static_cast<u32>(const_buffers.size() + global_buffers.size() +
+                                texel_buffers.size() + samplers.size() + images.size());
+    }
+
     std::vector<ConstBufferEntry> const_buffers;
     std::vector<GlobalBufferEntry> global_buffers;
+    std::vector<TexelBufferEntry> texel_buffers;
     std::vector<SamplerEntry> samplers;
+    std::vector<ImageEntry> images;
     std::set<u32> attributes;
     std::array<bool, Maxwell::NumClipDistances> clip_distances{};
     std::size_t shader_length{};
-    Sirit::Id entry_function{};
-    std::vector<Sirit::Id> interfaces;
+    bool uses_warps{};
 };
 
-using DecompilerResult = std::pair<std::unique_ptr<Sirit::Module>, ShaderEntries>;
+struct Specialization final {
+    u32 base_binding{};
 
-DecompilerResult Decompile(const VKDevice& device, const VideoCommon::Shader::ShaderIR& ir,
-                           Tegra::Engines::ShaderType stage);
+    // Compute specific
+    std::array<u32, 3> workgroup_size{};
+    u32 shared_memory_size{};
 
-} // namespace Vulkan::VKShader
+    // Graphics specific
+    Maxwell::PrimitiveTopology primitive_topology{};
+    std::optional<float> point_size{};
+    std::array<Maxwell::VertexAttribute::Type, Maxwell::NumVertexAttributes> attribute_types{};
+
+    // Tessellation specific
+    struct {
+        Maxwell::TessellationPrimitive primitive{};
+        Maxwell::TessellationSpacing spacing{};
+        bool clockwise{};
+    } tessellation;
+};
+// Old gcc versions don't consider this trivially copyable.
+// static_assert(std::is_trivially_copyable_v<Specialization>);
+
+struct SPIRVShader {
+    std::vector<u32> code;
+    ShaderEntries entries;
+};
+
+ShaderEntries GenerateShaderEntries(const VideoCommon::Shader::ShaderIR& ir);
+
+std::vector<u32> Decompile(const VKDevice& device, const VideoCommon::Shader::ShaderIR& ir,
+                           Tegra::Engines::ShaderType stage, const Specialization& specialization);
+
+} // namespace Vulkan

src/video_core/renderer_vulkan/vk_swapchain.cpp

@@ -19,12 +19,18 @@
 namespace Vulkan {
 
 namespace {
-vk::SurfaceFormatKHR ChooseSwapSurfaceFormat(const std::vector<vk::SurfaceFormatKHR>& formats) {
+
+vk::SurfaceFormatKHR ChooseSwapSurfaceFormat(const std::vector<vk::SurfaceFormatKHR>& formats,
+                                             bool srgb) {
     if (formats.size() == 1 && formats[0].format == vk::Format::eUndefined) {
-        return {vk::Format::eB8G8R8A8Unorm, vk::ColorSpaceKHR::eSrgbNonlinear};
+        vk::SurfaceFormatKHR format;
+        format.format = vk::Format::eB8G8R8A8Unorm;
+        format.colorSpace = vk::ColorSpaceKHR::eSrgbNonlinear;
+        return format;
     }
-    const auto& found = std::find_if(formats.begin(), formats.end(), [](const auto& format) {
-        return format.format == vk::Format::eB8G8R8A8Unorm &&
+    const auto& found = std::find_if(formats.begin(), formats.end(), [srgb](const auto& format) {
+        const auto request_format = srgb ? vk::Format::eB8G8R8A8Srgb : vk::Format::eB8G8R8A8Unorm;
+        return format.format == request_format &&
                format.colorSpace == vk::ColorSpaceKHR::eSrgbNonlinear;
     });
     return found != formats.end() ? *found : formats[0];
@@ -51,28 +57,26 @@ vk::Extent2D ChooseSwapExtent(const vk::SurfaceCapabilitiesKHR& capabilities, u3
                              std::min(capabilities.maxImageExtent.height, extent.height));
     return extent;
 }
-} // namespace
+
+} // Anonymous namespace
 
 VKSwapchain::VKSwapchain(vk::SurfaceKHR surface, const VKDevice& device)
     : surface{surface}, device{device} {}
 
 VKSwapchain::~VKSwapchain() = default;
 
-void VKSwapchain::Create(u32 width, u32 height) {
-    const auto dev = device.GetLogical();
+void VKSwapchain::Create(u32 width, u32 height, bool srgb) {
     const auto& dld = device.GetDispatchLoader();
     const auto physical_device = device.GetPhysical();
-
-    const vk::SurfaceCapabilitiesKHR capabilities{
-        physical_device.getSurfaceCapabilitiesKHR(surface, dld)};
+    const auto capabilities{physical_device.getSurfaceCapabilitiesKHR(surface, dld)};
     if (capabilities.maxImageExtent.width == 0 || capabilities.maxImageExtent.height == 0) {
         return;
     }
 
-    dev.waitIdle(dld);
+    device.GetLogical().waitIdle(dld);
     Destroy();
 
-    CreateSwapchain(capabilities, width, height);
+    CreateSwapchain(capabilities, width, height, srgb);
     CreateSemaphores();
     CreateImageViews();
 
@@ -107,7 +111,7 @@ bool VKSwapchain::Present(vk::Semaphore render_semaphore, VKFence& fence) {
         break;
     case vk::Result::eErrorOutOfDateKHR:
         if (current_width > 0 && current_height > 0) {
-            Create(current_width, current_height);
+            Create(current_width, current_height, current_srgb);
             recreated = true;
         }
         break;
@@ -129,23 +133,19 @@ bool VKSwapchain::HasFramebufferChanged(const Layout::FramebufferLayout& framebu
 }
 
 void VKSwapchain::CreateSwapchain(const vk::SurfaceCapabilitiesKHR& capabilities, u32 width,
-                                  u32 height) {
-    const auto dev{device.GetLogical()};
+                                  u32 height, bool srgb) {
     const auto& dld{device.GetDispatchLoader()};
     const auto physical_device{device.GetPhysical()};
+    const auto formats{physical_device.getSurfaceFormatsKHR(surface, dld)};
+    const auto present_modes{physical_device.getSurfacePresentModesKHR(surface, dld)};
 
-    const std::vector<vk::SurfaceFormatKHR> formats{
-        physical_device.getSurfaceFormatsKHR(surface, dld)};
-
-    const std::vector<vk::PresentModeKHR> present_modes{
-        physical_device.getSurfacePresentModesKHR(surface, dld)};
-
-    const vk::SurfaceFormatKHR surface_format{ChooseSwapSurfaceFormat(formats)};
+    const vk::SurfaceFormatKHR surface_format{ChooseSwapSurfaceFormat(formats, srgb)};
     const vk::PresentModeKHR present_mode{ChooseSwapPresentMode(present_modes)};
     extent = ChooseSwapExtent(capabilities, width, height);
 
     current_width = extent.width;
     current_height = extent.height;
+    current_srgb = srgb;
 
     u32 requested_image_count{capabilities.minImageCount + 1};
     if (capabilities.maxImageCount > 0 && requested_image_count > capabilities.maxImageCount) {
@@ -169,6 +169,7 @@ void VKSwapchain::CreateSwapchain(const vk::SurfaceCapabilitiesKHR& capabilities
         swapchain_ci.imageSharingMode = vk::SharingMode::eExclusive;
     }
 
+    const auto dev{device.GetLogical()};
     swapchain = dev.createSwapchainKHRUnique(swapchain_ci, nullptr, dld);
 
     images = dev.getSwapchainImagesKHR(*swapchain, dld);

src/video_core/renderer_vulkan/vk_swapchain.h

@@ -24,7 +24,7 @@ public:
     ~VKSwapchain();
 
     /// Creates (or recreates) the swapchain with a given size.
-    void Create(u32 width, u32 height);
+    void Create(u32 width, u32 height, bool srgb);
 
     /// Acquires the next image in the swapchain, waits as needed.
     void AcquireNextImage();
@@ -60,8 +60,13 @@ public:
         return image_format;
     }
 
+    bool GetSrgbState() const {
+        return current_srgb;
+    }
+
 private:
-    void CreateSwapchain(const vk::SurfaceCapabilitiesKHR& capabilities, u32 width, u32 height);
+    void CreateSwapchain(const vk::SurfaceCapabilitiesKHR& capabilities, u32 width, u32 height,
+                         bool srgb);
     void CreateSemaphores();
     void CreateImageViews();
 
@@ -87,6 +92,7 @@ private:
 
     u32 current_width{};
     u32 current_height{};
+    bool current_srgb{};
 };
 
 } // namespace Vulkan

src/video_core/shader/decode/arithmetic_integer.cpp

@@ -130,6 +130,25 @@ u32 ShaderIR::DecodeArithmeticInteger(NodeBlock& bb, u32 pc) {
         SetRegister(bb, instr.gpr0, value);
         break;
     }
+    case OpCode::Id::FLO_R:
+    case OpCode::Id::FLO_C:
+    case OpCode::Id::FLO_IMM: {
+        Node value;
+        if (instr.flo.invert) {
+            op_b = Operation(OperationCode::IBitwiseNot, NO_PRECISE, std::move(op_b));
+        }
+        if (instr.flo.is_signed) {
+            value = Operation(OperationCode::IBitMSB, NO_PRECISE, std::move(op_b));
+        } else {
+            value = Operation(OperationCode::UBitMSB, NO_PRECISE, std::move(op_b));
+        }
+        if (instr.flo.sh) {
+            value =
+                Operation(OperationCode::UBitwiseXor, NO_PRECISE, std::move(value), Immediate(31));
+        }
+        SetRegister(bb, instr.gpr0, std::move(value));
+        break;
+    }
     case OpCode::Id::SEL_C:
     case OpCode::Id::SEL_R:
     case OpCode::Id::SEL_IMM: {

src/video_core/shader/decode/conversion.cpp

@@ -63,12 +63,11 @@ u32 ShaderIR::DecodeConversion(NodeBlock& bb, u32 pc) {
     case OpCode::Id::I2F_R:
     case OpCode::Id::I2F_C:
     case OpCode::Id::I2F_IMM: {
-        UNIMPLEMENTED_IF(instr.conversion.int_src.selector != 0);
         UNIMPLEMENTED_IF(instr.conversion.dst_size == Register::Size::Long);
         UNIMPLEMENTED_IF_MSG(instr.generates_cc,
                              "Condition codes generation in I2F is not implemented");
 
-        Node value = [&]() {
+        Node value = [&] {
             switch (opcode->get().GetId()) {
             case OpCode::Id::I2F_R:
                 return GetRegister(instr.gpr20);
@@ -81,7 +80,19 @@ u32 ShaderIR::DecodeConversion(NodeBlock& bb, u32 pc) {
                 return Immediate(0);
             }
         }();
+
         const bool input_signed = instr.conversion.is_input_signed;
+
+        if (instr.conversion.src_size == Register::Size::Byte) {
+            const u32 offset = static_cast<u32>(instr.conversion.int_src.selector) * 8;
+            if (offset > 0) {
+                value = SignedOperation(OperationCode::ILogicalShiftRight, input_signed,
+                                        std::move(value), Immediate(offset));
+            }
+        } else {
+            UNIMPLEMENTED_IF(instr.conversion.int_src.selector != 0);
+        }
+
         value = ConvertIntegerSize(value, instr.conversion.src_size, input_signed);
         value = GetOperandAbsNegInteger(value, instr.conversion.abs_a, false, input_signed);
         value = SignedOperation(OperationCode::FCastInteger, input_signed, PRECISE, value);

src/video_core/shader/decode/memory.cpp

@@ -21,7 +21,24 @@ using Tegra::Shader::OpCode;
 using Tegra::Shader::Register;
 
 namespace {
-u32 GetUniformTypeElementsCount(Tegra::Shader::UniformType uniform_type) {
+
+u32 GetLdgMemorySize(Tegra::Shader::UniformType uniform_type) {
+    switch (uniform_type) {
+    case Tegra::Shader::UniformType::UnsignedByte:
+    case Tegra::Shader::UniformType::Single:
+        return 1;
+    case Tegra::Shader::UniformType::Double:
+        return 2;
+    case Tegra::Shader::UniformType::Quad:
+    case Tegra::Shader::UniformType::UnsignedQuad:
+        return 4;
+    default:
+        UNIMPLEMENTED_MSG("Unimplemented size={}!", static_cast<u32>(uniform_type));
+        return 1;
+    }
+}
+
+u32 GetStgMemorySize(Tegra::Shader::UniformType uniform_type) {
     switch (uniform_type) {
     case Tegra::Shader::UniformType::Single:
         return 1;
@@ -35,6 +52,7 @@ u32 GetUniformTypeElementsCount(Tegra::Shader::UniformType uniform_type) {
         return 1;
     }
 }
+
 } // Anonymous namespace
 
 u32 ShaderIR::DecodeMemory(NodeBlock& bb, u32 pc) {
@@ -168,7 +186,7 @@ u32 ShaderIR::DecodeMemory(NodeBlock& bb, u32 pc) {
         const auto [real_address_base, base_address, descriptor] =
             TrackGlobalMemory(bb, instr, false);
 
-        const u32 count = GetUniformTypeElementsCount(type);
+        const u32 count = GetLdgMemorySize(type);
         if (!real_address_base || !base_address) {
             // Tracking failed, load zeroes.
             for (u32 i = 0; i < count; ++i) {
@@ -179,12 +197,22 @@ u32 ShaderIR::DecodeMemory(NodeBlock& bb, u32 pc) {
 
         for (u32 i = 0; i < count; ++i) {
             const Node it_offset = Immediate(i * 4);
-            const Node real_address =
-                Operation(OperationCode::UAdd, NO_PRECISE, real_address_base, it_offset);
-            const Node gmem = MakeNode<GmemNode>(real_address, base_address, descriptor);
+            const Node real_address = Operation(OperationCode::UAdd, real_address_base, it_offset);
+            Node gmem = MakeNode<GmemNode>(real_address, base_address, descriptor);
+
+            if (type == Tegra::Shader::UniformType::UnsignedByte) {
+                // To handle unaligned loads get the byte used to dereferenced global memory
+                // and extract that byte from the loaded uint32.
+                Node byte = Operation(OperationCode::UBitwiseAnd, real_address, Immediate(3));
+                byte = Operation(OperationCode::ULogicalShiftLeft, std::move(byte), Immediate(3));
+
+                gmem = Operation(OperationCode::UBitfieldExtract, std::move(gmem), std::move(byte),
+                                 Immediate(8));
+            }
 
             SetTemporary(bb, i, gmem);
         }
+
         for (u32 i = 0; i < count; ++i) {
             SetRegister(bb, instr.gpr0.Value() + i, GetTemporary(i));
         }
@@ -196,28 +224,28 @@ u32 ShaderIR::DecodeMemory(NodeBlock& bb, u32 pc) {
         UNIMPLEMENTED_IF_MSG((instr.attribute.fmt20.immediate.Value() % sizeof(u32)) != 0,
                              "Unaligned attribute loads are not supported");
 
-        u64 next_element = instr.attribute.fmt20.element;
-        auto next_index = static_cast<u64>(instr.attribute.fmt20.index.Value());
+        u64 element = instr.attribute.fmt20.element;
+        auto index = static_cast<u64>(instr.attribute.fmt20.index.Value());
 
-        const auto StoreNextElement = [&](u32 reg_offset) {
-            const auto dest = GetOutputAttribute(static_cast<Attribute::Index>(next_index),
-                                                 next_element, GetRegister(instr.gpr39));
+        const u32 num_words = static_cast<u32>(instr.attribute.fmt20.size.Value()) + 1;
+        for (u32 reg_offset = 0; reg_offset < num_words; ++reg_offset) {
+            Node dest;
+            if (instr.attribute.fmt20.patch) {
+                const u32 offset = static_cast<u32>(index) * 4 + static_cast<u32>(element);
+                dest = MakeNode<PatchNode>(offset);
+            } else {
+                dest = GetOutputAttribute(static_cast<Attribute::Index>(index), element,
+                                          GetRegister(instr.gpr39));
+            }
             const auto src = GetRegister(instr.gpr0.Value() + reg_offset);
 
             bb.push_back(Operation(OperationCode::Assign, dest, src));
 
-            // Load the next attribute element into the following register. If the element
-            // to load goes beyond the vec4 size, load the first element of the next
-            // attribute.
-            next_element = (next_element + 1) % 4;
-            next_index = next_index + (next_element == 0 ? 1 : 0);
-        };
-
-        const u32 num_words = static_cast<u32>(instr.attribute.fmt20.size.Value()) + 1;
-        for (u32 reg_offset = 0; reg_offset < num_words; ++reg_offset) {
-            StoreNextElement(reg_offset);
+            // Load the next attribute element into the following register. If the element to load
+            // goes beyond the vec4 size, load the first element of the next attribute.
+            element = (element + 1) % 4;
+            index = index + (element == 0 ? 1 : 0);
         }
-
         break;
     }
     case OpCode::Id::ST_L:
@@ -274,7 +302,7 @@ u32 ShaderIR::DecodeMemory(NodeBlock& bb, u32 pc) {
             break;
         }
 
-        const u32 count = GetUniformTypeElementsCount(type);
+        const u32 count = GetStgMemorySize(type);
         for (u32 i = 0; i < count; ++i) {
             const Node it_offset = Immediate(i * 4);
             const Node real_address = Operation(OperationCode::UAdd, real_address_base, it_offset);

src/video_core/shader/decode/other.cpp

@@ -69,6 +69,8 @@ u32 ShaderIR::DecodeOther(NodeBlock& bb, u32 pc) {
     case OpCode::Id::MOV_SYS: {
         const Node value = [this, instr] {
             switch (instr.sys20) {
+            case SystemVariable::InvocationId:
+                return Operation(OperationCode::InvocationId);
             case SystemVariable::Ydirection:
                 return Operation(OperationCode::YNegate);
             case SystemVariable::InvocationInfo:
@@ -255,6 +257,12 @@ u32 ShaderIR::DecodeOther(NodeBlock& bb, u32 pc) {
         SetRegister(bb, instr.gpr0, GetRegister(instr.gpr8));
         break;
     }
+    case OpCode::Id::MEMBAR: {
+        UNIMPLEMENTED_IF(instr.membar.type != Tegra::Shader::MembarType::GL);
+        UNIMPLEMENTED_IF(instr.membar.unknown != Tegra::Shader::MembarUnknown::Default);
+        bb.push_back(Operation(OperationCode::MemoryBarrierGL));
+        break;
+    }
     case OpCode::Id::DEPBAR: {
         LOG_DEBUG(HW_GPU, "DEPBAR instruction is stubbed");
         break;

src/video_core/shader/decode/texture.cpp

@@ -107,8 +107,8 @@ u32 ShaderIR::DecodeTexture(NodeBlock& bb, u32 pc) {
         break;
     }
     case OpCode::Id::TLD4S: {
-        UNIMPLEMENTED_IF_MSG(instr.tld4s.UsesMiscMode(TextureMiscMode::AOFFI),
-                             "AOFFI is not implemented");
+        const bool uses_aoffi = instr.tld4s.UsesMiscMode(TextureMiscMode::AOFFI);
+        UNIMPLEMENTED_IF_MSG(uses_aoffi, "AOFFI is not implemented");
 
         const bool depth_compare = instr.tld4s.UsesMiscMode(TextureMiscMode::DC);
         const Node op_a = GetRegister(instr.gpr8);
@@ -116,29 +116,86 @@ u32 ShaderIR::DecodeTexture(NodeBlock& bb, u32 pc) {
 
         // TODO(Subv): Figure out how the sampler type is encoded in the TLD4S instruction.
         std::vector<Node> coords;
+        Node dc_reg;
         if (depth_compare) {
             // Note: TLD4S coordinate encoding works just like TEXS's
             const Node op_y = GetRegister(instr.gpr8.Value() + 1);
             coords.push_back(op_a);
             coords.push_back(op_y);
-            coords.push_back(op_b);
+            dc_reg = uses_aoffi ? GetRegister(instr.gpr20.Value() + 1) : op_b;
         } else {
             coords.push_back(op_a);
-            coords.push_back(op_b);
+            if (uses_aoffi) {
+                const Node op_y = GetRegister(instr.gpr8.Value() + 1);
+                coords.push_back(op_y);
+            } else {
+                coords.push_back(op_b);
+            }
+            dc_reg = {};
         }
         const Node component = Immediate(static_cast<u32>(instr.tld4s.component));
 
         const SamplerInfo info{TextureType::Texture2D, false, depth_compare};
-        const auto& sampler = GetSampler(instr.sampler, info);
+        const Sampler& sampler = *GetSampler(instr.sampler, info);
 
         Node4 values;
         for (u32 element = 0; element < values.size(); ++element) {
             auto coords_copy = coords;
-            MetaTexture meta{sampler, {}, {}, {}, {}, {}, component, element};
+            MetaTexture meta{sampler, {}, dc_reg, {}, {}, {}, {}, component, element};
             values[element] = Operation(OperationCode::TextureGather, meta, std::move(coords_copy));
         }
 
-        WriteTexsInstructionFloat(bb, instr, values, true);
+        if (instr.tld4s.fp16_flag) {
+            WriteTexsInstructionHalfFloat(bb, instr, values, true);
+        } else {
+            WriteTexsInstructionFloat(bb, instr, values, true);
+        }
+        break;
+    }
+    case OpCode::Id::TXD_B:
+        is_bindless = true;
+        [[fallthrough]];
+    case OpCode::Id::TXD: {
+        UNIMPLEMENTED_IF_MSG(instr.txd.UsesMiscMode(TextureMiscMode::AOFFI),
+                             "AOFFI is not implemented");
+        UNIMPLEMENTED_IF_MSG(instr.txd.is_array != 0, "TXD Array is not implemented");
+
+        u64 base_reg = instr.gpr8.Value();
+        const auto derivate_reg = instr.gpr20.Value();
+        const auto texture_type = instr.txd.texture_type.Value();
+        const auto coord_count = GetCoordCount(texture_type);
+
+        const Sampler* sampler = is_bindless
+                                     ? GetBindlessSampler(base_reg, {{texture_type, false, false}})
+                                     : GetSampler(instr.sampler, {{texture_type, false, false}});
+        Node4 values;
+        if (sampler == nullptr) {
+            for (u32 element = 0; element < values.size(); ++element) {
+                values[element] = Immediate(0);
+            }
+            WriteTexInstructionFloat(bb, instr, values);
+            break;
+        }
+        if (is_bindless) {
+            base_reg++;
+        }
+
+        std::vector<Node> coords;
+        std::vector<Node> derivates;
+        for (std::size_t i = 0; i < coord_count; ++i) {
+            coords.push_back(GetRegister(base_reg + i));
+            const std::size_t derivate = i * 2;
+            derivates.push_back(GetRegister(derivate_reg + derivate));
+            derivates.push_back(GetRegister(derivate_reg + derivate + 1));
+        }
+
+        for (u32 element = 0; element < values.size(); ++element) {
+            MetaTexture meta{*sampler, {}, {}, {}, derivates, {}, {}, {}, element};
+            values[element] = Operation(OperationCode::TextureGradient, std::move(meta), coords);
+        }
+
+        WriteTexInstructionFloat(bb, instr, values);
+
         break;
     }
     case OpCode::Id::TXQ_B:
@@ -148,9 +205,24 @@ u32 ShaderIR::DecodeTexture(NodeBlock& bb, u32 pc) {
         // TODO: The new commits on the texture refactor, change the way samplers work.
         // Sadly, not all texture instructions specify the type of texture their sampler
         // uses. This must be fixed at a later instance.
-        const auto& sampler =
+        const Sampler* sampler =
             is_bindless ? GetBindlessSampler(instr.gpr8) : GetSampler(instr.sampler);
 
+        if (sampler == nullptr) {
+            u32 indexer = 0;
+            for (u32 element = 0; element < 4; ++element) {
+                if (!instr.txq.IsComponentEnabled(element)) {
+                    continue;
+                }
+                const Node value = Immediate(0);
+                SetTemporary(bb, indexer++, value);
+            }
+            for (u32 i = 0; i < indexer; ++i) {
+                SetRegister(bb, instr.gpr0.Value() + i, GetTemporary(i));
+            }
+            break;
+        }
+
         u32 indexer = 0;
         switch (instr.txq.query_type) {
         case Tegra::Shader::TextureQueryType::Dimension: {
@@ -158,7 +230,7 @@ u32 ShaderIR::DecodeTexture(NodeBlock& bb, u32 pc) {
                 if (!instr.txq.IsComponentEnabled(element)) {
                     continue;
                 }
-                MetaTexture meta{sampler, {}, {}, {}, {}, {}, {}, element};
+                MetaTexture meta{*sampler, {}, {}, {}, {}, {}, {}, {}, element};
                 const Node value =
                     Operation(OperationCode::TextureQueryDimensions, meta,
                               GetRegister(instr.gpr8.Value() + (is_bindless ? 1 : 0)));
@@ -184,9 +256,24 @@ u32 ShaderIR::DecodeTexture(NodeBlock& bb, u32 pc) {
 
         auto texture_type = instr.tmml.texture_type.Value();
         const bool is_array = instr.tmml.array != 0;
-        const auto& sampler =
+        const Sampler* sampler =
             is_bindless ? GetBindlessSampler(instr.gpr20) : GetSampler(instr.sampler);
 
+        if (sampler == nullptr) {
+            u32 indexer = 0;
+            for (u32 element = 0; element < 2; ++element) {
+                if (!instr.tmml.IsComponentEnabled(element)) {
+                    continue;
+                }
+                const Node value = Immediate(0);
+                SetTemporary(bb, indexer++, value);
+            }
+            for (u32 i = 0; i < indexer; ++i) {
+                SetRegister(bb, instr.gpr0.Value() + i, GetTemporary(i));
+            }
+            break;
+        }
+
         std::vector<Node> coords;
 
         // TODO: Add coordinates for different samplers once other texture types are implemented.
@@ -212,7 +299,7 @@ u32 ShaderIR::DecodeTexture(NodeBlock& bb, u32 pc) {
                 continue;
             }
             auto params = coords;
-            MetaTexture meta{sampler, {}, {}, {}, {}, {}, {}, element};
+            MetaTexture meta{*sampler, {}, {}, {}, {}, {}, {}, {}, element};
             const Node value = Operation(OperationCode::TextureQueryLod, meta, std::move(params));
             SetTemporary(bb, indexer++, value);
         }
@@ -268,7 +355,7 @@ ShaderIR::SamplerInfo ShaderIR::GetSamplerInfo(std::optional<SamplerInfo> sample
                        sampler->is_buffer != 0};
 }
 
-const Sampler& ShaderIR::GetSampler(const Tegra::Shader::Sampler& sampler,
+const Sampler* ShaderIR::GetSampler(const Tegra::Shader::Sampler& sampler,
                                     std::optional<SamplerInfo> sampler_info) {
     const auto offset = static_cast<u32>(sampler.index.Value());
     const auto info = GetSamplerInfo(sampler_info, offset);
@@ -280,21 +367,24 @@ const Sampler& ShaderIR::GetSampler(const Tegra::Shader::Sampler& sampler,
     if (it != used_samplers.end()) {
         ASSERT(!it->IsBindless() && it->GetType() == info.type && it->IsArray() == info.is_array &&
                it->IsShadow() == info.is_shadow && it->IsBuffer() == info.is_buffer);
-        return *it;
+        return &(*it);
     }
 
     // Otherwise create a new mapping for this sampler
     const auto next_index = static_cast<u32>(used_samplers.size());
-    return used_samplers.emplace_back(next_index, offset, info.type, info.is_array, info.is_shadow,
-                                      info.is_buffer);
+    return &used_samplers.emplace_back(next_index, offset, info.type, info.is_array, info.is_shadow,
+                                       info.is_buffer);
 }
 
-const Sampler& ShaderIR::GetBindlessSampler(Tegra::Shader::Register reg,
+const Sampler* ShaderIR::GetBindlessSampler(Tegra::Shader::Register reg,
                                             std::optional<SamplerInfo> sampler_info) {
     const Node sampler_register = GetRegister(reg);
     const auto [base_sampler, buffer, offset] =
         TrackCbuf(sampler_register, global_code, static_cast<s64>(global_code.size()));
     ASSERT(base_sampler != nullptr);
+    if (base_sampler == nullptr) {
+        return nullptr;
+    }
 
     const auto info = GetSamplerInfo(sampler_info, offset, buffer);
 
@@ -307,13 +397,13 @@ const Sampler& ShaderIR::GetBindlessSampler(Tegra::Shader::Register reg,
     if (it != used_samplers.end()) {
         ASSERT(it->IsBindless() && it->GetType() == info.type && it->IsArray() == info.is_array &&
                it->IsShadow() == info.is_shadow);
-        return *it;
+        return &(*it);
     }
 
     // Otherwise create a new mapping for this sampler
     const auto next_index = static_cast<u32>(used_samplers.size());
-    return used_samplers.emplace_back(next_index, offset, buffer, info.type, info.is_array,
-                                      info.is_shadow, info.is_buffer);
+    return &used_samplers.emplace_back(next_index, offset, buffer, info.type, info.is_array,
+                                       info.is_shadow, info.is_buffer);
 }
 
 void ShaderIR::WriteTexInstructionFloat(NodeBlock& bb, Instruction instr, const Node4& components) {
@@ -356,14 +446,14 @@ void ShaderIR::WriteTexsInstructionFloat(NodeBlock& bb, Instruction instr, const
 }
 
 void ShaderIR::WriteTexsInstructionHalfFloat(NodeBlock& bb, Instruction instr,
-                                             const Node4& components) {
+                                             const Node4& components, bool ignore_mask) {
     // TEXS.F16 destionation registers are packed in two registers in pairs (just like any half
     // float instruction).
 
     Node4 values;
     u32 dest_elem = 0;
     for (u32 component = 0; component < 4; ++component) {
-        if (!instr.texs.IsComponentEnabled(component))
+        if (!instr.texs.IsComponentEnabled(component) && !ignore_mask)
             continue;
         values[dest_elem++] = components[component];
     }
@@ -399,8 +489,15 @@ Node4 ShaderIR::GetTextureCode(Instruction instr, TextureType texture_type,
                          "This method is not supported.");
 
     const SamplerInfo info{texture_type, is_array, is_shadow, false};
-    const auto& sampler =
+    const Sampler* sampler =
         is_bindless ? GetBindlessSampler(*bindless_reg, info) : GetSampler(instr.sampler, info);
+    Node4 values;
+    if (sampler == nullptr) {
+        for (u32 element = 0; element < values.size(); ++element) {
+            values[element] = Immediate(0);
+        }
+        return values;
+    }
 
     const bool lod_needed = process_mode == TextureProcessMode::LZ ||
                             process_mode == TextureProcessMode::LL ||
@@ -439,10 +536,9 @@ Node4 ShaderIR::GetTextureCode(Instruction instr, TextureType texture_type,
         }
     }
 
-    Node4 values;
     for (u32 element = 0; element < values.size(); ++element) {
         auto copy_coords = coords;
-        MetaTexture meta{sampler, array, depth_compare, aoffi, bias, lod, {}, element};
+        MetaTexture meta{*sampler, array, depth_compare, aoffi, {}, bias, lod, {}, element};
         values[element] = Operation(read_method, meta, std::move(copy_coords));
     }
 
@@ -555,8 +651,15 @@ Node4 ShaderIR::GetTld4Code(Instruction instr, TextureType texture_type, bool de
     u64 parameter_register = instr.gpr20.Value();
 
     const SamplerInfo info{texture_type, is_array, depth_compare, false};
-    const auto& sampler = is_bindless ? GetBindlessSampler(parameter_register++, info)
-                                      : GetSampler(instr.sampler, info);
+    const Sampler* sampler = is_bindless ? GetBindlessSampler(parameter_register++, info)
+                                         : GetSampler(instr.sampler, info);
+    Node4 values;
+    if (sampler == nullptr) {
+        for (u32 element = 0; element < values.size(); ++element) {
+            values[element] = Immediate(0);
+        }
+        return values;
+    }
 
     std::vector<Node> aoffi;
     if (is_aoffi) {
@@ -571,10 +674,9 @@ Node4 ShaderIR::GetTld4Code(Instruction instr, TextureType texture_type, bool de
     const Node component = is_bindless ? Immediate(static_cast<u32>(instr.tld4_b.component))
                                        : Immediate(static_cast<u32>(instr.tld4.component));
 
-    Node4 values;
     for (u32 element = 0; element < values.size(); ++element) {
         auto coords_copy = coords;
-        MetaTexture meta{sampler, GetRegister(array_register), dc, aoffi, {}, {}, component,
+        MetaTexture meta{*sampler, GetRegister(array_register), dc, aoffi, {}, {}, {}, component,
                          element};
         values[element] = Operation(OperationCode::TextureGather, meta, std::move(coords_copy));
     }
@@ -603,12 +705,12 @@ Node4 ShaderIR::GetTldCode(Tegra::Shader::Instruction instr) {
     // const Node aoffi_register{is_aoffi ? GetRegister(gpr20_cursor++) : nullptr};
     // const Node multisample{is_multisample ? GetRegister(gpr20_cursor++) : nullptr};
 
-    const auto& sampler = GetSampler(instr.sampler);
+    const auto& sampler = *GetSampler(instr.sampler);
 
     Node4 values;
     for (u32 element = 0; element < values.size(); ++element) {
         auto coords_copy = coords;
-        MetaTexture meta{sampler, array_register, {}, {}, {}, lod, {}, element};
+        MetaTexture meta{sampler, array_register, {}, {}, {}, {}, lod, {}, element};
         values[element] = Operation(OperationCode::TexelFetch, meta, std::move(coords_copy));
     }
 
@@ -616,7 +718,7 @@ Node4 ShaderIR::GetTldCode(Tegra::Shader::Instruction instr) {
 }
 
 Node4 ShaderIR::GetTldsCode(Instruction instr, TextureType texture_type, bool is_array) {
-    const auto& sampler = GetSampler(instr.sampler);
+    const Sampler& sampler = *GetSampler(instr.sampler);
 
     const std::size_t type_coord_count = GetCoordCount(texture_type);
     const bool lod_enabled = instr.tlds.GetTextureProcessMode() == TextureProcessMode::LL;
@@ -641,19 +743,24 @@ Node4 ShaderIR::GetTldsCode(Instruction instr, TextureType texture_type, bool is
     // When lod is used always is in gpr20
     const Node lod = lod_enabled ? GetRegister(instr.gpr20) : Immediate(0);
 
-    // Fill empty entries from the guest sampler.
+    // Fill empty entries from the guest sampler
     const std::size_t entry_coord_count = GetCoordCount(sampler.GetType());
     if (type_coord_count != entry_coord_count) {
         LOG_WARNING(HW_GPU, "Bound and built texture types mismatch");
-    }
-    for (std::size_t i = type_coord_count; i < entry_coord_count; ++i) {
-        coords.push_back(GetRegister(Register::ZeroIndex));
+
+        // When the size is higher we insert zeroes
+        for (std::size_t i = type_coord_count; i < entry_coord_count; ++i) {
+            coords.push_back(GetRegister(Register::ZeroIndex));
+        }
+
+        // Then we ensure the size matches the number of entries (dropping unused values)
+        coords.resize(entry_coord_count);
     }
 
     Node4 values;
     for (u32 element = 0; element < values.size(); ++element) {
         auto coords_copy = coords;
-        MetaTexture meta{sampler, array, {}, {}, {}, lod, {}, element};
+        MetaTexture meta{sampler, array, {}, {}, {}, {}, lod, {}, element};
         values[element] = Operation(OperationCode::TexelFetch, meta, std::move(coords_copy));
     }
     return values;

src/video_core/shader/decode/warp.cpp

@@ -38,6 +38,9 @@ u32 ShaderIR::DecodeWarp(NodeBlock& bb, u32 pc) {
     const Instruction instr = {program_code[pc]};
     const auto opcode = OpCode::Decode(instr);
 
+    // Signal the backend that this shader uses warp instructions.
+    uses_warps = true;
+
     switch (opcode->get().GetId()) {
     case OpCode::Id::VOTE: {
         const Node value = GetPredicate(instr.vote.value, instr.vote.negate_value != 0);

src/video_core/shader/node.h

@@ -68,6 +68,7 @@ enum class OperationCode {
     IBitfieldInsert,       /// (MetaArithmetic, int base, int insert, int offset, int bits) -> int
     IBitfieldExtract,      /// (MetaArithmetic, int value, int offset, int offset) -> int
     IBitCount,             /// (MetaArithmetic, int) -> int
+    IBitMSB,               /// (MetaArithmetic, int) -> int
 
     UAdd,                  /// (MetaArithmetic, uint a, uint b) -> uint
     UMul,                  /// (MetaArithmetic, uint a, uint b) -> uint
@@ -86,6 +87,7 @@ enum class OperationCode {
     UBitfieldInsert,  /// (MetaArithmetic, uint base, uint insert, int offset, int bits) -> uint
     UBitfieldExtract, /// (MetaArithmetic, uint value, int offset, int offset) -> uint
     UBitCount,        /// (MetaArithmetic, uint) -> uint
+    UBitMSB,          /// (MetaArithmetic, uint) -> uint
 
     HAdd,       /// (MetaArithmetic, f16vec2 a, f16vec2 b) -> f16vec2
     HMul,       /// (MetaArithmetic, f16vec2 a, f16vec2 b) -> f16vec2
@@ -149,6 +151,7 @@ enum class OperationCode {
     TextureQueryDimensions, /// (MetaTexture, float a) -> float4
     TextureQueryLod,        /// (MetaTexture, float[N] coords) -> float4
     TexelFetch,             /// (MetaTexture, int[N], int) -> float4
+    TextureGradient,        /// (MetaTexture, float[N] coords, float[N*2] derivates) -> float4
 
     ImageLoad,  /// (MetaImage, int[N] coords) -> void
     ImageStore, /// (MetaImage, int[N] coords) -> void
@@ -169,6 +172,7 @@ enum class OperationCode {
     EmitVertex,   /// () -> void
     EndPrimitive, /// () -> void
 
+    InvocationId,       /// () -> int
     YNegate,            /// () -> float
     LocalInvocationIdX, /// () -> uint
     LocalInvocationIdY, /// () -> uint
@@ -185,6 +189,8 @@ enum class OperationCode {
     ThreadId,       /// () -> uint
     ShuffleIndexed, /// (uint value, uint index) -> uint
 
+    MemoryBarrierGL, /// () -> void
+
     Amount,
 };
 
@@ -210,13 +216,14 @@ class PredicateNode;
 class AbufNode;
 class CbufNode;
 class LmemNode;
+class PatchNode;
 class SmemNode;
 class GmemNode;
 class CommentNode;
 
-using NodeData =
-    std::variant<OperationNode, ConditionalNode, GprNode, ImmediateNode, InternalFlagNode,
-                 PredicateNode, AbufNode, CbufNode, LmemNode, SmemNode, GmemNode, CommentNode>;
+using NodeData = std::variant<OperationNode, ConditionalNode, GprNode, ImmediateNode,
+                              InternalFlagNode, PredicateNode, AbufNode, PatchNode, CbufNode,
+                              LmemNode, SmemNode, GmemNode, CommentNode>;
 using Node = std::shared_ptr<NodeData>;
 using Node4 = std::array<Node, 4>;
 using NodeBlock = std::vector<Node>;
@@ -367,6 +374,7 @@ struct MetaTexture {
     Node array;
     Node depth_compare;
     std::vector<Node> aoffi;
+    std::vector<Node> derivates;
     Node bias;
     Node lod;
     Node component{};
@@ -538,6 +546,19 @@ private:
     u32 element{};
 };
 
+/// Patch memory (used to communicate tessellation stages).
+class PatchNode final {
+public:
+    explicit PatchNode(u32 offset) : offset{offset} {}
+
+    u32 GetOffset() const {
+        return offset;
+    }
+
+private:
+    u32 offset{};
+};
+
 /// Constant buffer node, usually mapped to uniform buffers in GLSL
 class CbufNode final {
 public:

src/video_core/shader/shader_ir.h

@@ -137,6 +137,10 @@ public:
         return uses_vertex_id;
     }
 
+    bool UsesWarps() const {
+        return uses_warps;
+    }
+
     bool HasPhysicalAttributes() const {
         return uses_physical_attributes;
     }
@@ -309,11 +313,11 @@ private:
                                std::optional<u32> buffer = std::nullopt);
 
     /// Accesses a texture sampler
-    const Sampler& GetSampler(const Tegra::Shader::Sampler& sampler,
+    const Sampler* GetSampler(const Tegra::Shader::Sampler& sampler,
                               std::optional<SamplerInfo> sampler_info = std::nullopt);
 
     /// Accesses a texture sampler for a bindless texture.
-    const Sampler& GetBindlessSampler(Tegra::Shader::Register reg,
+    const Sampler* GetBindlessSampler(Tegra::Shader::Register reg,
                                       std::optional<SamplerInfo> sampler_info = std::nullopt);
 
     /// Accesses an image.
@@ -334,7 +338,7 @@ private:
     void WriteTexsInstructionFloat(NodeBlock& bb, Tegra::Shader::Instruction instr,
                                    const Node4& components, bool ignore_mask = false);
     void WriteTexsInstructionHalfFloat(NodeBlock& bb, Tegra::Shader::Instruction instr,
-                                       const Node4& components);
+                                       const Node4& components, bool ignore_mask = false);
 
     Node4 GetTexCode(Tegra::Shader::Instruction instr, Tegra::Shader::TextureType texture_type,
                      Tegra::Shader::TextureProcessMode process_mode, bool depth_compare,
@@ -415,6 +419,7 @@ private:
     bool uses_physical_attributes{}; // Shader uses AL2P or physical attribute read/writes
     bool uses_instance_id{};
     bool uses_vertex_id{};
+    bool uses_warps{};
 
     Tegra::Shader::Header header;
 };

src/video_core/shader/track.cpp

@@ -7,6 +7,7 @@
 #include <variant>
 
 #include "common/common_types.h"
+#include "video_core/shader/node.h"
 #include "video_core/shader/shader_ir.h"
 
 namespace VideoCommon::Shader {

src/video_core/texture_cache/surface_params.cpp

@@ -246,6 +246,16 @@ SurfaceParams SurfaceParams::CreateForFermiCopySurface(
     return params;
 }
 
+VideoCore::Surface::SurfaceTarget SurfaceParams::ExpectedTarget(
+    const VideoCommon::Shader::Sampler& entry) {
+    return TextureTypeToSurfaceTarget(entry.GetType(), entry.IsArray());
+}
+
+VideoCore::Surface::SurfaceTarget SurfaceParams::ExpectedTarget(
+    const VideoCommon::Shader::Image& entry) {
+    return ImageTypeToSurfaceTarget(entry.GetType());
+}
+
 bool SurfaceParams::IsLayered() const {
     switch (target) {
     case SurfaceTarget::Texture1DArray:
@@ -382,4 +392,42 @@ std::string SurfaceParams::TargetName() const {
     }
 }
 
+u32 SurfaceParams::GetBlockSize() const {
+    const u32 x = 64U << block_width;
+    const u32 y = 8U << block_height;
+    const u32 z = 1U << block_depth;
+    return x * y * z;
+}
+
+std::pair<u32, u32> SurfaceParams::GetBlockXY() const {
+    const u32 x_pixels = 64U / GetBytesPerPixel();
+    const u32 x = x_pixels << block_width;
+    const u32 y = 8U << block_height;
+    return {x, y};
+}
+
+std::tuple<u32, u32, u32> SurfaceParams::GetBlockOffsetXYZ(u32 offset) const {
+    const auto div_ceil = [](const u32 x, const u32 y) { return ((x + y - 1) / y); };
+    const u32 block_size = GetBlockSize();
+    const u32 block_index = offset / block_size;
+    const u32 gob_offset = offset % block_size;
+    const u32 gob_index = gob_offset / static_cast<u32>(Tegra::Texture::GetGOBSize());
+    const u32 x_gob_pixels = 64U / GetBytesPerPixel();
+    const u32 x_block_pixels = x_gob_pixels << block_width;
+    const u32 y_block_pixels = 8U << block_height;
+    const u32 z_block_pixels = 1U << block_depth;
+    const u32 x_blocks = div_ceil(width, x_block_pixels);
+    const u32 y_blocks = div_ceil(height, y_block_pixels);
+    const u32 z_blocks = div_ceil(depth, z_block_pixels);
+    const u32 base_x = block_index % x_blocks;
+    const u32 base_y = (block_index / x_blocks) % y_blocks;
+    const u32 base_z = (block_index / (x_blocks * y_blocks)) % z_blocks;
+    u32 x = base_x * x_block_pixels;
+    u32 y = base_y * y_block_pixels;
+    u32 z = base_z * z_block_pixels;
+    z += gob_index >> block_height;
+    y += (gob_index * 8U) % y_block_pixels;
+    return {x, y, z};
+}
+
 } // namespace VideoCommon

src/video_core/texture_cache/surface_params.h

@@ -4,6 +4,8 @@
 
 #pragma once
 
+#include <utility>
+
 #include "common/alignment.h"
 #include "common/bit_util.h"
 #include "common/cityhash.h"
@@ -45,6 +47,14 @@ public:
     static SurfaceParams CreateForFermiCopySurface(
         const Tegra::Engines::Fermi2D::Regs::Surface& config);
 
+    /// Obtains the texture target from a shader's sampler entry.
+    static VideoCore::Surface::SurfaceTarget ExpectedTarget(
+        const VideoCommon::Shader::Sampler& entry);
+
+    /// Obtains the texture target from a shader's sampler entry.
+    static VideoCore::Surface::SurfaceTarget ExpectedTarget(
+        const VideoCommon::Shader::Image& entry);
+
     std::size_t Hash() const {
         return static_cast<std::size_t>(
             Common::CityHash64(reinterpret_cast<const char*>(this), sizeof(*this)));
@@ -128,6 +138,15 @@ public:
 
     std::size_t GetConvertedMipmapSize(u32 level) const;
 
+    /// Get this texture Tegra Block size in guest memory layout
+    u32 GetBlockSize() const;
+
+    /// Get X, Y coordinates max sizes of a single block.
+    std::pair<u32, u32> GetBlockXY() const;
+
+    /// Get the offset in x, y, z coordinates from a memory offset
+    std::tuple<u32, u32, u32> GetBlockOffsetXYZ(u32 offset) const;
+
     /// Returns the size of a layer in bytes in guest memory.
     std::size_t GetGuestLayerSize() const {
         return GetLayerSize(false, false);
@@ -261,7 +280,8 @@ private:
 
     /// Returns the size of all mipmap levels and aligns as needed.
     std::size_t GetInnerMemorySize(bool as_host_size, bool layer_only, bool uncompressed) const {
-        return GetLayerSize(as_host_size, uncompressed) * (layer_only ? 1U : depth);
+        return GetLayerSize(as_host_size, uncompressed) *
+               (layer_only ? 1U : (is_layered ? depth : 1U));
     }
 
     /// Returns the size of a layer

src/video_core/texture_cache/texture_cache.h

@@ -95,10 +95,16 @@ public:
         std::lock_guard lock{mutex};
         const auto gpu_addr{tic.Address()};
         if (!gpu_addr) {
-            return {};
+            return GetNullSurface(SurfaceParams::ExpectedTarget(entry));
+        }
+
+        const auto host_ptr{system.GPU().MemoryManager().GetPointer(gpu_addr)};
+        const auto cache_addr{ToCacheAddr(host_ptr)};
+        if (!cache_addr) {
+            return GetNullSurface(SurfaceParams::ExpectedTarget(entry));
         }
         const auto params{SurfaceParams::CreateForTexture(format_lookup_table, tic, entry)};
-        const auto [surface, view] = GetSurface(gpu_addr, params, true, false);
+        const auto [surface, view] = GetSurface(gpu_addr, cache_addr, params, true, false);
         if (guard_samplers) {
             sampled_textures.push_back(surface);
         }
@@ -110,10 +116,15 @@ public:
         std::lock_guard lock{mutex};
         const auto gpu_addr{tic.Address()};
         if (!gpu_addr) {
-            return {};
+            return GetNullSurface(SurfaceParams::ExpectedTarget(entry));
+        }
+        const auto host_ptr{system.GPU().MemoryManager().GetPointer(gpu_addr)};
+        const auto cache_addr{ToCacheAddr(host_ptr)};
+        if (!cache_addr) {
+            return GetNullSurface(SurfaceParams::ExpectedTarget(entry));
         }
         const auto params{SurfaceParams::CreateForImage(format_lookup_table, tic, entry)};
-        const auto [surface, view] = GetSurface(gpu_addr, params, true, false);
+        const auto [surface, view] = GetSurface(gpu_addr, cache_addr, params, true, false);
         if (guard_samplers) {
             sampled_textures.push_back(surface);
         }
@@ -143,11 +154,17 @@ public:
             SetEmptyDepthBuffer();
             return {};
         }
+        const auto host_ptr{system.GPU().MemoryManager().GetPointer(gpu_addr)};
+        const auto cache_addr{ToCacheAddr(host_ptr)};
+        if (!cache_addr) {
+            SetEmptyDepthBuffer();
+            return {};
+        }
         const auto depth_params{SurfaceParams::CreateForDepthBuffer(
             system, regs.zeta_width, regs.zeta_height, regs.zeta.format,
             regs.zeta.memory_layout.block_width, regs.zeta.memory_layout.block_height,
             regs.zeta.memory_layout.block_depth, regs.zeta.memory_layout.type)};
-        auto surface_view = GetSurface(gpu_addr, depth_params, preserve_contents, true);
+        auto surface_view = GetSurface(gpu_addr, cache_addr, depth_params, preserve_contents, true);
         if (depth_buffer.target)
             depth_buffer.target->MarkAsRenderTarget(false, NO_RT);
         depth_buffer.target = surface_view.first;
@@ -180,8 +197,16 @@ public:
             return {};
         }
 
-        auto surface_view = GetSurface(gpu_addr, SurfaceParams::CreateForFramebuffer(system, index),
-                                       preserve_contents, true);
+        const auto host_ptr{system.GPU().MemoryManager().GetPointer(gpu_addr)};
+        const auto cache_addr{ToCacheAddr(host_ptr)};
+        if (!cache_addr) {
+            SetEmptyColorBuffer(index);
+            return {};
+        }
+
+        auto surface_view =
+            GetSurface(gpu_addr, cache_addr, SurfaceParams::CreateForFramebuffer(system, index),
+                       preserve_contents, true);
         if (render_targets[index].target)
             render_targets[index].target->MarkAsRenderTarget(false, NO_RT);
         render_targets[index].target = surface_view.first;
@@ -230,8 +255,14 @@ public:
         const GPUVAddr src_gpu_addr = src_config.Address();
         const GPUVAddr dst_gpu_addr = dst_config.Address();
         DeduceBestBlit(src_params, dst_params, src_gpu_addr, dst_gpu_addr);
-        std::pair<TSurface, TView> dst_surface = GetSurface(dst_gpu_addr, dst_params, true, false);
-        std::pair<TSurface, TView> src_surface = GetSurface(src_gpu_addr, src_params, true, false);
+        const auto dst_host_ptr{system.GPU().MemoryManager().GetPointer(dst_gpu_addr)};
+        const auto dst_cache_addr{ToCacheAddr(dst_host_ptr)};
+        const auto src_host_ptr{system.GPU().MemoryManager().GetPointer(src_gpu_addr)};
+        const auto src_cache_addr{ToCacheAddr(src_host_ptr)};
+        std::pair<TSurface, TView> dst_surface =
+            GetSurface(dst_gpu_addr, dst_cache_addr, dst_params, true, false);
+        std::pair<TSurface, TView> src_surface =
+            GetSurface(src_gpu_addr, src_cache_addr, src_params, true, false);
         ImageBlit(src_surface.second, dst_surface.second, copy_config);
         dst_surface.first->MarkAsModified(true, Tick());
     }
@@ -347,13 +378,6 @@ protected:
         return new_surface;
     }
 
-    std::pair<TSurface, TView> GetFermiSurface(
-        const Tegra::Engines::Fermi2D::Regs::Surface& config) {
-        SurfaceParams params = SurfaceParams::CreateForFermiCopySurface(config);
-        const GPUVAddr gpu_addr = config.Address();
-        return GetSurface(gpu_addr, params, true, false);
-    }
-
     Core::System& system;
 
 private:
@@ -591,6 +615,86 @@ private:
         return {{new_surface, new_surface->GetMainView()}};
     }
 
+    /**
+     * Takes care of managing 3D textures and its slices. Does HLE methods for reconstructing the 3D
+     * textures within the GPU if possible. Falls back to LLE when it isn't possible to use any of
+     * the HLE methods.
+     *
+     * @param overlaps          The overlapping surfaces registered in the cache.
+     * @param params            The parameters on the new surface.
+     * @param gpu_addr          The starting address of the new surface.
+     * @param cache_addr        The starting address of the new surface on physical memory.
+     * @param preserve_contents Indicates that the new surface should be loaded from memory or
+     *                          left blank.
+     */
+    std::optional<std::pair<TSurface, TView>> Manage3DSurfaces(std::vector<TSurface>& overlaps,
+                                                               const SurfaceParams& params,
+                                                               const GPUVAddr gpu_addr,
+                                                               const CacheAddr cache_addr,
+                                                               bool preserve_contents) {
+        if (params.target == SurfaceTarget::Texture3D) {
+            bool failed = false;
+            if (params.num_levels > 1) {
+                // We can't handle mipmaps in 3D textures yet, better fallback to LLE approach
+                return std::nullopt;
+            }
+            TSurface new_surface = GetUncachedSurface(gpu_addr, params);
+            bool modified = false;
+            for (auto& surface : overlaps) {
+                const SurfaceParams& src_params = surface->GetSurfaceParams();
+                if (src_params.target != SurfaceTarget::Texture2D) {
+                    failed = true;
+                    break;
+                }
+                if (src_params.height != params.height) {
+                    failed = true;
+                    break;
+                }
+                if (src_params.block_depth != params.block_depth ||
+                    src_params.block_height != params.block_height) {
+                    failed = true;
+                    break;
+                }
+                const u32 offset = static_cast<u32>(surface->GetCacheAddr() - cache_addr);
+                const auto [x, y, z] = params.GetBlockOffsetXYZ(offset);
+                modified |= surface->IsModified();
+                const CopyParams copy_params(0, 0, 0, 0, 0, z, 0, 0, params.width, params.height,
+                                             1);
+                ImageCopy(surface, new_surface, copy_params);
+            }
+            if (failed) {
+                return std::nullopt;
+            }
+            for (const auto& surface : overlaps) {
+                Unregister(surface);
+            }
+            new_surface->MarkAsModified(modified, Tick());
+            Register(new_surface);
+            auto view = new_surface->GetMainView();
+            return {{std::move(new_surface), view}};
+        } else {
+            for (const auto& surface : overlaps) {
+                if (!surface->MatchTarget(params.target)) {
+                    if (overlaps.size() == 1 && surface->GetCacheAddr() == cache_addr) {
+                        if (Settings::values.use_accurate_gpu_emulation) {
+                            return std::nullopt;
+                        }
+                        Unregister(surface);
+                        return InitializeSurface(gpu_addr, params, preserve_contents);
+                    }
+                    return std::nullopt;
+                }
+                if (surface->GetCacheAddr() != cache_addr) {
+                    continue;
+                }
+                if (surface->MatchesStructure(params) == MatchStructureResult::FullMatch) {
+                    return {{surface, surface->GetMainView()}};
+                }
+            }
+            return InitializeSurface(gpu_addr, params, preserve_contents);
+        }
+    }
+
     /**
      * Gets the starting address and parameters of a candidate surface and tries
      * to find a matching surface within the cache. This is done in 3 big steps:
@@ -614,22 +718,9 @@ private:
      *                          left blank.
      * @param is_render         Whether or not the surface is a render target.
      **/
-    std::pair<TSurface, TView> GetSurface(const GPUVAddr gpu_addr, const SurfaceParams& params,
-                                          bool preserve_contents, bool is_render) {
-        const auto host_ptr{system.GPU().MemoryManager().GetPointer(gpu_addr)};
-        const auto cache_addr{ToCacheAddr(host_ptr)};
-
-        // Step 0: guarantee a valid surface
-        if (!cache_addr) {
-            // Return a null surface if it's invalid
-            SurfaceParams new_params = params;
-            new_params.width = 1;
-            new_params.height = 1;
-            new_params.depth = 1;
-            new_params.block_height = 0;
-            new_params.block_depth = 0;
-            return InitializeSurface(gpu_addr, new_params, false);
-        }
+    std::pair<TSurface, TView> GetSurface(const GPUVAddr gpu_addr, const CacheAddr cache_addr,
+                                          const SurfaceParams& params, bool preserve_contents,
+                                          bool is_render) {
 
         // Step 1
         // Check Level 1 Cache for a fast structural match. If candidate surface
@@ -676,6 +767,15 @@ private:
             }
         }
 
+        // Check if it's a 3D texture
+        if (params.block_depth > 0) {
+            auto surface =
+                Manage3DSurfaces(overlaps, params, gpu_addr, cache_addr, preserve_contents);
+            if (surface) {
+                return *surface;
+            }
+        }
+
         // Split cases between 1 overlap or many.
         if (overlaps.size() == 1) {
             TSurface current_surface = overlaps[0];
@@ -793,6 +893,41 @@ private:
         }
     }
 
+    /**
+     * Gets a null surface based on a target texture.
+     * @param target The target of the null surface.
+     */
+    TView GetNullSurface(SurfaceTarget target) {
+        const u32 i_target = static_cast<u32>(target);
+        if (const auto it = invalid_cache.find(i_target); it != invalid_cache.end()) {
+            return it->second->GetMainView();
+        }
+        SurfaceParams params{};
+        params.target = target;
+        params.is_tiled = false;
+        params.srgb_conversion = false;
+        params.is_layered = false;
+        params.block_width = 0;
+        params.block_height = 0;
+        params.block_depth = 0;
+        params.tile_width_spacing = 1;
+        params.width = 1;
+        params.height = 1;
+        params.depth = 1;
+        params.pitch = 4;
+        params.num_levels = 1;
+        params.emulated_levels = 1;
+        params.pixel_format = VideoCore::Surface::PixelFormat::RGBA16F;
+        params.type = VideoCore::Surface::SurfaceType::ColorTexture;
+        auto surface = CreateSurface(0ULL, params);
+        invalid_memory.clear();
+        invalid_memory.resize(surface->GetHostSizeInBytes(), 0U);
+        surface->UploadTexture(invalid_memory);
+        surface->MarkAsModified(false, Tick());
+        invalid_cache.emplace(i_target, surface);
+        return surface->GetMainView();
+    }
+
     /**
      * Gets the a source and destination starting address and parameters,
      * and tries to deduce if they are supposed to be depth textures. If so, their
@@ -991,6 +1126,11 @@ private:
 
     std::vector<TSurface> sampled_textures;
 
+    /// This cache stores null surfaces in order to be used as a placeholder
+    /// for invalid texture calls.
+    std::unordered_map<u32, TSurface> invalid_cache;
+    std::vector<u8> invalid_memory;
+
     StagingCache staging_cache;
     std::recursive_mutex mutex;
 };

src/video_core/textures/decoders.h

@@ -12,6 +12,10 @@ namespace Tegra::Texture {
 
 // GOBSize constant. Calculated by 64 bytes in x multiplied by 8 y coords, represents
 // an small rect of (64/bytes_per_pixel)X8.
+inline std::size_t GetGOBSize() {
+    return 512;
+}
+
 inline std::size_t GetGOBSizeShift() {
     return 9;
 }