From f32b430326af86bd956f3d49d1bf7b6a1b9c5e31 Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Tue, 27 Oct 2026 19:30:20 -0400 Subject: Switch 2: native memory/images, GPU-accessible attribute, TSG subcontexts, syncpoint shim - gsv 111/112/114/423/429/430/314/315/435 track native memory and images with real ids (meanings from changes.patch); 1510/435 use its block-linear layout. - svc 0x80 sets a real GpuAccessible (0x20) memory attribute (from changes.patch). - Characteristics report TSG subcontexts (flag 22, 64 per TSG at 0xb4); NvRm otherwise allows one subcontext and a second NVN queue fails its setup. - Syncpoint shim: flag 26, address-space ioctl 0xB, and GPU reads of the shim page return guest syncpoint values (scheme from changes.patch). - EDEN_NX2_IPC_TRACE also logs TIPC commands; gdbrsp.py watch/mem improvements. --- src/video_core/memory_manager.cpp | 13 +++++++++++++ 1 file changed, 13 insertions(+) (limited to 'src/video_core/memory_manager.cpp') diff --git a/src/video_core/memory_manager.cpp b/src/video_core/memory_manager.cpp index 8ffdcb6276..0290500e7c 100644 --- a/src/video_core/memory_manager.cpp +++ b/src/video_core/memory_manager.cpp @@ -56,6 +56,10 @@ MemoryManager::MemoryManager(Core::System& system_, u64 address_space_bits_, GPU : MemoryManager(system_, system_.Host1x().MemoryManager(), address_space_bits_, split_address_, big_page_bits_, page_bits_) {} +void MemoryManager::SetNativeSyncpointShim(GPUVAddr address) { + native_syncpoint_shim.store(address, std::memory_order_release); +} + MemoryManager::~MemoryManager() = default; MemoryManager::EntryType MemoryManager::GetEntry(size_t position, bool is_big_page) const { @@ -231,6 +235,15 @@ std::optional MemoryManager::GpuToCpuAddress(GPUVAddr addr, std::size_t s template T MemoryManager::Read(GPUVAddr addr) const { + if constexpr (sizeof(T) == sizeof(u32)) { + const auto shim = native_syncpoint_shim.load(std::memory_order_acquire); + if (shim && addr >= shim && addr - shim < NumShimSyncpoints * sizeof(u32) && + Common::IsAligned(addr - shim, sizeof(u32))) { + return static_cast(system.Host1x().GetSyncpointManager().GetGuestSyncpointValue( + static_cast((addr - shim) / sizeof(u32)))); + } + } + if (auto page_pointer{GetPointer(addr)}; page_pointer) { // NOTE: Avoid adding any extra logic to this fast-path block T value; -- cgit v1.2.3 From bf1de8f40130af5d3b490ee40dcc2fb20d460a14 Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Tue, 27 Oct 2026 20:53:15 -0400 Subject: Switch 2: suspend channels on semaphore acquires; fix pushbuffer desyncs NVN2 gates command buffers with semaphores the CPU releases after writing them. Eden read past unmet SEMAPHORED acquires, so it executed stale commands and desynchronised. Channels now suspend on an unmet acquire and resume from the same word once it is met (scheduler keeps the waiting channels, the GPU thread polls them), as in changes.patch. Also: methods split across GPFIFO entries around a called command buffer are set aside and resumed; subchannel 2 defaults to inline-to-memory; unbound subchannels and out-of-range compute/DMA methods are dropped instead of crashing; 4-byte semaphore releases no longer write 16 bytes; Ampere SEM_PAYLOAD_HI/SEM_EXECUTE no longer spin as Kepler acquires; unmapped GPU reads log instead of asserting. Debug aids: EDEN_NX2_PB_CHECK (pushbuffer changed after submit), EDEN_NX2_QUERY_TRACE, a one-time pushbuffer dump. --- src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp | 12 ++ src/video_core/control/channel_state.cpp | 6 +- src/video_core/control/scheduler.cpp | 32 +++++ src/video_core/control/scheduler.h | 8 ++ src/video_core/dma_pusher.cpp | 157 ++++++++++++++++++++-- src/video_core/dma_pusher.h | 49 ++++++- src/video_core/engines/kepler_compute.cpp | 15 ++- src/video_core/engines/maxwell_dma.cpp | 6 +- src/video_core/engines/puller.cpp | 133 +++++++++++++++++- src/video_core/engines/puller.h | 23 ++++ src/video_core/gpu_thread.cpp | 14 +- src/video_core/memory_manager.cpp | 5 +- src/video_core/nx2.h | 35 ++++- src/video_core/query_cache/query_cache.h | 5 + 14 files changed, 473 insertions(+), 27 deletions(-) (limited to 'src/video_core/memory_manager.cpp') diff --git a/src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp b/src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp index a72ccd9b4f..cf3948c321 100644 --- a/src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp +++ b/src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp @@ -4,7 +4,9 @@ // SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include #include +#include #include "common/assert.h" #include "common/logging.h" #include "core/core.h" @@ -422,6 +424,16 @@ NvResult nvhost_gpu::SubmitGPFIFOImpl(IoctlSubmitGpfifo& params, Tegra::CommandL LOG_DEBUG(Service_NVDRV, "nx2 GPFIFO entry raw={:#018x} addr={:#x} size={:#x}", entry.raw, entry.addr.Value(), entry.size.Value()); } + static const bool check = std::getenv("EDEN_NX2_PB_CHECK") != nullptr; + if (check) { + std::vector words; + for (const auto& entry : entries.command_lists) { + words.resize(entry.size); + system.ApplicationMemory().ReadBlock(entry.addr.Value(), words.data(), + words.size() * sizeof(u32)); + entries.submit_hashes.push_back(Tegra::HashPushbufferWords(words)); + } + } } auto& flags = params.flags; diff --git a/src/video_core/control/channel_state.cpp b/src/video_core/control/channel_state.cpp index 12a452dfc0..a1a8110510 100644 --- a/src/video_core/control/channel_state.cpp +++ b/src/video_core/control/channel_state.cpp @@ -21,6 +21,7 @@ namespace { // Match NVK/Nouveau's initial pushbuffer subchannel layout. constexpr u32 Nvk3DSubchannel = 0; constexpr u32 NvkComputeSubchannel = 1; +constexpr u32 NvkM2mfSubchannel = 2; constexpr u32 Nvk2DSubchannel = 3; constexpr u32 NvkCopySubchannel = 4; @@ -29,7 +30,10 @@ void BindNvkDefaultSubchannels(ChannelState::Payload& payload) { dma_pusher.BindSubchannel(&payload.maxwell_3d, Nvk3DSubchannel, Engines::EngineTypes::Maxwell3D); dma_pusher.BindSubchannel(&payload.kepler_compute, NvkComputeSubchannel, Engines::EngineTypes::KeplerCompute); - // Subchannel 2 is M2MF there; Eden does not expose a 0x9039 engine yet. + // Subchannel 2 is M2MF/inline-to-memory there (NVN uses it without binding a class); + // KeplerMemory has the same LINE_LENGTH_IN/OFFSET_OUT/LAUNCH_DMA/LOAD_INLINE_DATA layout. + dma_pusher.BindSubchannel(&payload.kepler_memory, NvkM2mfSubchannel, + Engines::EngineTypes::KeplerMemory); dma_pusher.BindSubchannel(&payload.fermi_2d, Nvk2DSubchannel, Engines::EngineTypes::Fermi2D); dma_pusher.BindSubchannel(&payload.maxwell_dma, NvkCopySubchannel, Engines::EngineTypes::MaxwellDMA); diff --git a/src/video_core/control/scheduler.cpp b/src/video_core/control/scheduler.cpp index 8af13e415c..24f1db1f7c 100644 --- a/src/video_core/control/scheduler.cpp +++ b/src/video_core/control/scheduler.cpp @@ -4,7 +4,9 @@ // SPDX-FileCopyrightText: 2021 yuzu Emulator Project // SPDX-License-Identifier: GPL-3.0-or-later +#include #include +#include #include "common/assert.h" #include "video_core/control/channel_state.h" @@ -26,6 +28,36 @@ void Scheduler::Push(GPU& gpu, s32 channel, CommandList&& entries) { // Multiple channels can prepare their commands in parallel. channel_state->payload->dma_pusher.Push(std::move(entries)); channel_state->payload->dma_pusher.DispatchCalls(); + if (channel_state->payload->dma_pusher.HasPendingAcquire() && + std::find(pending_channels.begin(), pending_channels.end(), channel_state) == + pending_channels.end()) { + pending_channels.push_back(channel_state); + } + if (!pending_channels.empty()) { + ResumePending(gpu); + } +} + +void Scheduler::ResumePending(GPU& gpu) { + // Running one channel can release the semaphore another waits for: retry until none moves. + for (bool progress = true; progress;) { + progress = false; + for (const auto& state : pending_channels) { + auto& dma_pusher = state->payload->dma_pusher; + if (!dma_pusher.HasPendingAcquire() || !dma_pusher.CanResumeAcquire()) { + continue; + } + { + std::scoped_lock lk{scheduling_guard}; + gpu.BindChannel(state->bind_id); + } + dma_pusher.DispatchCalls(); + progress = true; + } + } + std::erase_if(pending_channels, [](const auto& state) { + return !state->payload->dma_pusher.HasPendingAcquire(); + }); } void Scheduler::DeclareChannel(std::shared_ptr new_channel) { diff --git a/src/video_core/control/scheduler.h b/src/video_core/control/scheduler.h index 342bea5025..3a3dcc58cf 100644 --- a/src/video_core/control/scheduler.h +++ b/src/video_core/control/scheduler.h @@ -8,6 +8,7 @@ #include #include +#include #include "common/container/unordered_map.h" #include "video_core/dma_pusher.h" @@ -24,11 +25,18 @@ class Scheduler { public: void Push(GPU& gpu, s32 channel, CommandList&& entries); + /// Runs channels whose pending semaphore acquire is now met. + void ResumePending(GPU& gpu); + [[nodiscard]] bool HasPendingChannels() const { + return !pending_channels.empty(); + } + void DeclareChannel(std::shared_ptr new_channel); private: ::Common::unordered_map> channels; std::mutex scheduling_guard; + std::vector> pending_channels; ///< Suspended on an acquire }; } // namespace Control diff --git a/src/video_core/dma_pusher.cpp b/src/video_core/dma_pusher.cpp index 2c0ba007ef..0889964337 100644 --- a/src/video_core/dma_pusher.cpp +++ b/src/video_core/dma_pusher.cpp @@ -4,13 +4,17 @@ // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include + #include "common/settings.h" #include "core/core.h" +#include "core/memory.h" #include "video_core/dma_pusher.h" #include "video_core/engines/maxwell_3d.h" #include "video_core/gpu.h" #include "video_core/guest_memory.h" #include "video_core/memory_manager.h" +#include "video_core/nx2.h" #include "video_core/rasterizer_interface.h" namespace Tegra { @@ -29,7 +33,7 @@ DmaPusher::DmaPusher(Core::System& system_, MemoryManager& memory_manager_, Cont DmaPusher::~DmaPusher() = default; void DmaPusher::DispatchCalls() { - dma_pushbuffer_subindex = 0; + // The sub-index is kept: a channel suspended on a semaphore resumes inside its command list. dma_state.is_last_call = true; while (system.IsPoweredOn()) { if (!Step()) { @@ -44,6 +48,9 @@ bool DmaPusher::Step() { if (!ib_enable || dma_pushbuffer.empty()) { return false; } + if (!puller.ResumeAcquire(*this)) { + return false; + } CommandList& command_list = dma_pushbuffer.front(); @@ -57,7 +64,13 @@ bool DmaPusher::Step() { } if (prefetch_size > 0) { - ProcessCommands(command_list.prefetch_command_list); + const std::span commands{command_list.prefetch_command_list}; + command_resume_offset += + ProcessCommands(commands.subspan(command_resume_offset), command_resume_offset); + if (puller.HasPendingAcquire()) { + return false; + } + command_resume_offset = 0; dma_pushbuffer.pop(); return true; } @@ -83,13 +96,33 @@ bool DmaPusher::Step() { subchannels[dma_state.subchannel]->current_dirty = memory_manager.IsMemoryDirty(dma_state.dma_get, words * sizeof(u32)); } } + if (command_resume_offset == 0) { + if (g_nx2_ampere) { + ResolveSplitMethod(header.addr); + } + if (dma_pushbuffer_subindex < command_list.submit_hashes.size()) { + CheckSubmitHash(header, command_list.submit_hashes[dma_pushbuffer_subindex]); + } + } + // A resumed entry is read again: the words after the acquire may be new by now. const bool use_safe = Settings::IsDMALevelDefault() ? Settings::IsGPULevelHigh() : Settings::IsDMALevelSafe(); if (use_safe) { Tegra::Memory::GpuGuestMemoryheaders(memory_manager, dma_state.dma_get, header.size, &command_headers); - ProcessCommands(headers); + const std::span commands{headers.data(), headers.size()}; + command_resume_offset += + ProcessCommands(commands.subspan(command_resume_offset), command_resume_offset); } else { Tegra::Memory::GpuGuestMemoryheaders(memory_manager, dma_state.dma_get, header.size, &command_headers); - ProcessCommands(headers); + const std::span commands{headers.data(), headers.size()}; + command_resume_offset += + ProcessCommands(commands.subspan(command_resume_offset), command_resume_offset); + } + if (puller.HasPendingAcquire()) { + return false; + } + command_resume_offset = 0; + if (dma_state.method_count != 0) { + split_method_end = header.addr + header.size * sizeof(u32); } } @@ -111,20 +144,23 @@ bool DmaPusher::Step() { return true; } -void DmaPusher::ProcessCommands(std::span commands) { +size_t DmaPusher::ProcessCommands(std::span commands, size_t word_offset) { for (size_t index = 0; index < commands.size();) { // Data word of methods command if (dma_state.method_count && dma_state.non_incrementing) { auto const& command_header = commands[index]; //must ref (MUltiMethod re) - dma_state.dma_word_offset = u32(index * sizeof(u32)); - const u32 max_write = u32(std::min(index + dma_state.method_count, commands.size()) - index); + dma_state.dma_word_offset = u64((word_offset + index) * sizeof(u32)); + u32 max_write = u32(std::min(index + dma_state.method_count, commands.size()) - index); + if (g_nx2_ampere && dma_state.method < non_puller_methods) { + max_write = 1; // a semaphore acquire may suspend the channel after any word + } CallMultiMethod(&command_header.argument, max_write); dma_state.method_count -= max_write; dma_state.is_last_call = true; index += max_write; } else if (dma_state.method_count) { auto const command_header = commands[index]; //can copy - dma_state.dma_word_offset = u32(index * sizeof(u32)); + dma_state.dma_word_offset = u64((word_offset + index) * sizeof(u32)); dma_state.is_last_call = dma_state.method_count <= 1; CallMethod(command_header.argument); dma_state.method += !dma_state.non_incrementing ? 1 : 0; @@ -163,6 +199,46 @@ void DmaPusher::ProcessCommands(std::span commands) { } index++; } + if (puller.HasPendingAcquire()) { + return index; + } + } + return commands.size(); +} + +// NVN on Switch 2 ends GPFIFO entries in the middle of a method's data and calls another command +// buffer (e.g. a compute dispatch) before the entry that carries the rest, which starts right +// where the first one ended. An entry starting anywhere else is such a call: run it with a clean +// method state and resume the unfinished method when the stream continues at that address. +void DmaPusher::ResolveSplitMethod(GPUVAddr entry_address) { + if (dma_state.method_count != 0 && entry_address != split_method_end) { + if (parked_method) { + static std::atomic reported{}; + if (reported.fetch_add(1) < 8) { + LOG_WARNING(HW_GPU, + "Dropping parked method {:#x} (sub {}, {} words left, resume {:#x})", + parked_method->method, parked_method->subchannel, + parked_method->method_count, parked_method->resume_address); + } + } + parked_method = ParkedMethod{ + .method = dma_state.method, + .subchannel = dma_state.subchannel, + .method_count = dma_state.method_count, + .non_incrementing = dma_state.non_incrementing, + .increment_once = dma_increment_once, + .resume_address = split_method_end, + }; + dma_state.method_count = 0; + dma_increment_once = false; + } else if (dma_state.method_count == 0 && parked_method && + entry_address == parked_method->resume_address) { + dma_state.method = parked_method->method; + dma_state.subchannel = parked_method->subchannel; + dma_state.method_count = parked_method->method_count; + dma_state.non_incrementing = parked_method->non_incrementing; + dma_increment_once = parked_method->increment_once; + parked_method.reset(); } } @@ -182,6 +258,10 @@ void DmaPusher::CallMethod(u32 argument) { }); } else { auto subchannel = subchannels[dma_state.subchannel]; + if (subchannel == nullptr) { + ReportUnboundSubchannel(); + return; + } if (!subchannel->execution_mask[dma_state.method]) { subchannel->method_sink.emplace_back(dma_state.method, argument); } else { @@ -197,12 +277,73 @@ void DmaPusher::CallMultiMethod(const u32* base_start, u32 num_methods) { puller.CallMultiMethod(*this, dma_state.method, dma_state.subchannel, base_start, num_methods, dma_state.method_count); } else { auto subchannel = subchannels[dma_state.subchannel]; + if (subchannel == nullptr) { + ReportUnboundSubchannel(); + return; + } subchannel->ConsumeSink(system); subchannel->current_dma_segment = dma_state.dma_get + dma_state.dma_word_offset; subchannel->CallMultiMethod(system, dma_state.method, base_start, num_methods, dma_state.method_count); } } +void DmaPusher::ReportUnboundSubchannel() { + // Methods sent to a subchannel with no engine (no SetObject yet) are dropped. + static std::atomic reported{}; + if (reported.fetch_add(1) < 8) { + LOG_WARNING(HW_GPU, "Dropping method {:#x} on unbound subchannel {}", dma_state.method, + dma_state.subchannel); + } +} + +void DmaPusher::CheckSubmitHash(const CommandListHeader& header, u64 submit_hash) { + std::vector words(header.size); + memory_manager.ReadBlockUnsafe(header.addr, words.data(), words.size() * sizeof(u32)); + static std::atomic checked{}; + static std::atomic changed{}; + ++checked; + if (HashPushbufferWords(words) != submit_hash) { + if (changed.fetch_add(1) < 16) { + LOG_WARNING(HW_GPU, "pushbuffer entry {:#x}+{:#x} changed between submit and execution", + header.addr.Value(), header.size.Value()); + } + } + if ((checked % 20000) == 0) { + LOG_WARNING(HW_GPU, "pushbuffer check: {} of {} entries changed after submit", + changed.load(), checked.load()); + } +} + +void DmaPusher::DumpCurrentEntry(const char* reason) { + LOG_WARNING(HW_GPU, "pushbuffer dump ({}): method={:#x} sub={} count={} non_inc={} get={:#x}", + reason, dma_state.method, dma_state.subchannel, dma_state.method_count, + dma_state.non_incrementing, dma_state.dma_get); + if (dma_pushbuffer.empty()) { + return; + } + const auto& lists = dma_pushbuffer.front().command_lists; + const size_t first = dma_pushbuffer_subindex >= 2 ? dma_pushbuffer_subindex - 2 : 0; + for (size_t i = first; i <= dma_pushbuffer_subindex && i < lists.size(); ++i) { + const CommandListHeader& header = lists[i]; + LOG_WARNING(HW_GPU, " entry {} of {}: addr={:#x} size={:#x}", i, lists.size(), + header.addr.Value(), header.size.Value()); + for (u64 word = 0; word < std::min(header.size, 0x80); word += 8) { + std::string gpu_line; + std::string cpu_line; + for (u64 j = word; j < std::min(header.size, word + 8); ++j) { + const GPUVAddr address = header.addr + j * sizeof(u32); + const u32 gpu_value = memory_manager.Read(address); + const u32 cpu_value = system.ApplicationMemory().Read32(address); + gpu_line += fmt::format(" {:08x}", gpu_value); + cpu_line += gpu_value == cpu_value ? std::string(" ........") + : fmt::format(" {:08x}", cpu_value); + } + LOG_WARNING(HW_GPU, " +{:03x}:{}", word * 4, gpu_line); + LOG_WARNING(HW_GPU, " cpu :{}", cpu_line); + } + } +} + void DmaPusher::BindRasterizer(VideoCore::RasterizerInterface* rasterizer_) { rasterizer = rasterizer_; } diff --git a/src/video_core/dma_pusher.h b/src/video_core/dma_pusher.h index 62af628511..e735839d30 100644 --- a/src/video_core/dma_pusher.h +++ b/src/video_core/dma_pusher.h @@ -8,6 +8,7 @@ #include #include +#include #include #include #include @@ -114,8 +115,19 @@ struct CommandList final { boost::container::small_vector command_lists; boost::container::small_vector prefetch_command_list; + /// EDEN_NX2_PB_CHECK: hash of each entry's words when it was submitted. + std::vector submit_hashes; }; +/// FNV-1a over pushbuffer words (EDEN_NX2_PB_CHECK). +inline u64 HashPushbufferWords(std::span words) { + u64 hash = 0xcbf29ce484222325ULL; + for (const u32 word : words) { + hash = (hash ^ word) * 0x100000001b3ULL; + } + return hash; +} + /// @brief The DmaPusher class implements DMA submission to FIFOs, providing an area of memory that the /// emulated app fills with commands and tells PFIFO to process. The pushbuffers are then assembled /// into a "command stream" consisting of 32-bit words that make up "commands". @@ -132,6 +144,14 @@ public: void DispatchCalls(); + /// Switch 2: true while the channel is suspended on a semaphore acquire. + [[nodiscard]] bool HasPendingAcquire() const { + return puller.HasPendingAcquire(); + } + [[nodiscard]] bool CanResumeAcquire() { + return puller.CanResumeAcquire(*this); + } + void BindSubchannel(Engines::EngineInterface* engine, u32 subchannel_id, Engines::EngineTypes engine_type) { subchannels[subchannel_id] = engine; subchannel_type[subchannel_id] = engine_type; @@ -143,12 +163,24 @@ private: static constexpr u32 non_puller_methods = 0x40; static constexpr u32 max_subchannels = 8; bool Step(); - void ProcessCommands(std::span commands); + /// Returns how many words were processed: fewer than given when a semaphore acquire + /// suspended the channel. + size_t ProcessCommands(std::span commands, size_t word_offset = 0); + size_t command_resume_offset{}; ///< Words of the current entry done before a suspension void SetState(const CommandHeader& command_header); void CallMethod(u32 argument); void CallMultiMethod(const u32* base_start, u32 num_methods); + void ReportUnboundSubchannel(); + +public: + /// Logs the GPFIFO entry being processed (and the one before it), read through the GPU MMU + /// and straight from guest memory, to find pushbuffer desyncs. Debug aid for Switch 2. + void DumpCurrentEntry(const char* reason); + void CheckSubmitHash(const CommandListHeader& header, u64 submit_hash); + +private: public: Common::ScratchBuffer command_headers; ///< Buffer for list of commands fetched at once @@ -172,6 +204,21 @@ public: Control::ChannelState& channel_state; DmaState dma_state{}; + + // Switch 2: a method whose data continues in a later GPFIFO entry, set aside while another + // command buffer called in between runs. + struct ParkedMethod { + u32 method; + u32 subchannel; + u32 method_count; + bool non_incrementing; + bool increment_once; + GPUVAddr resume_address; ///< Where the entry that continues the method starts + }; + void ResolveSplitMethod(GPUVAddr entry_address); + std::optional parked_method; + GPUVAddr split_method_end{}; ///< End of the last entry that ended inside a method + std::array subchannels{}; std::array subchannel_type; diff --git a/src/video_core/engines/kepler_compute.cpp b/src/video_core/engines/kepler_compute.cpp index 96a12539a2..166bbd6d53 100644 --- a/src/video_core/engines/kepler_compute.cpp +++ b/src/video_core/engines/kepler_compute.cpp @@ -35,13 +35,24 @@ void KeplerCompute::BindRasterizer(VideoCore::RasterizerInterface* rasterizer_) void KeplerCompute::ConsumeSinkImpl(Core::System& system) { for (auto [method, value] : method_sink) { - regs.reg_array[method] = value; + if (method < Regs::NUM_REGS) { + regs.reg_array[method] = value; + } } method_sink.clear(); } void KeplerCompute::CallMethod(Core::System& system, u32 method, u32 method_argument, bool is_last_call) { - ASSERT_MSG(method < Regs::NUM_REGS, "Invalid KeplerCompute register, increase the size of the Regs structure"); + if (method >= Regs::NUM_REGS) { + // Ampere compute (Switch 2) uses methods past Kepler's register file; writing them would + // run off reg_array, so drop them. + static bool reported = false; + if (!reported) { + reported = true; + LOG_WARNING(HW_GPU, "Ignoring KeplerCompute method {:#x} past the register file", method); + } + return; + } regs.reg_array[method] = method_argument; diff --git a/src/video_core/engines/maxwell_dma.cpp b/src/video_core/engines/maxwell_dma.cpp index ed4675ed12..5ab953e510 100644 --- a/src/video_core/engines/maxwell_dma.cpp +++ b/src/video_core/engines/maxwell_dma.cpp @@ -42,7 +42,11 @@ void MaxwellDMA::ConsumeSinkImpl(Core::System& system) { } void MaxwellDMA::CallMethod(Core::System& system, u32 method, u32 method_argument, bool is_last_call) { - ASSERT_MSG(method < NUM_REGS, "Invalid MaxwellDMA register"); + if (method >= NUM_REGS) { + // Ampere copy (Switch 2) can use methods past this register file; don't write past it. + LOG_DEBUG(HW_GPU, "Ignoring MaxwellDMA method {:#x} past the register file", method); + return; + } regs.reg_array[method] = method_argument; diff --git a/src/video_core/engines/puller.cpp b/src/video_core/engines/puller.cpp index 9e78d586d1..97a58903fe 100644 --- a/src/video_core/engines/puller.cpp +++ b/src/video_core/engines/puller.cpp @@ -4,6 +4,7 @@ // SPDX-FileCopyrightText: 2022 yuzu Emulator Project // SPDX-License-Identifier: GPL-3.0-or-later +#include #include #include "common/assert.h" @@ -30,8 +31,58 @@ namespace { constexpr u32 Gf100BindClassMask = 0xffff; constexpr u32 Gf100BindValidMask = 0x1f0000 | Gf100BindClassMask; +// Reads a semaphore word, flushing a GPU-written copy out of the buffer cache first. +u32 ReadSemaphoreWord(DmaPusher& dma_pusher, GPUVAddr address) { + if (const auto device = dma_pusher.memory_manager.GpuToCpuAddress(address); + device && dma_pusher.rasterizer->MustFlushRegion(*device, sizeof(u32), + VideoCommon::CacheType::BufferCache)) { + dma_pusher.rasterizer->FlushRegion(*device, sizeof(u32), + VideoCommon::CacheType::BufferCache); + } + return dma_pusher.memory_manager.Read(address); +} + } // Anonymous namespace +bool Puller::PendingAcquire::Matches(u32 current) const { + switch (static_cast(operation)) { + case GpuSemaphoreOperation::AcquireEqual: + return current == value; + case GpuSemaphoreOperation::AcquireGequal: + return current >= value; + case GpuSemaphoreOperation::AcquireMask: + return (current & value) != 0; + default: + return true; + } +} + +bool Puller::CanResumeAcquire(DmaPusher& dma_pusher) { + if (!pending_acquire) { + return true; + } + const u32 current = ReadSemaphoreWord(dma_pusher, pending_acquire->address); + if (pending_acquire->Matches(current)) { + return true; + } + if (!pending_acquire->reported && + std::chrono::steady_clock::now() - pending_acquire->since > std::chrono::seconds{2}) { + pending_acquire->reported = true; + LOG_WARNING(HW_GPU, "Channel {} waits on semaphore {:#x} (op {}, {:#x}, now {:#x}) for 2 s", + dma_pusher.channel_state.bind_id, pending_acquire->address, + pending_acquire->operation, pending_acquire->value, current); + } + return false; +} + +bool Puller::ResumeAcquire(DmaPusher& dma_pusher) { + if (!CanResumeAcquire(dma_pusher)) { + return false; + } + pending_acquire.reset(); + return true; +} + void Puller::ProcessBindMethod(DmaPusher& dma_pusher, const MethodCall& method_call) { LOG_DEBUG(HW_GPU, "Binding subchannel {} to engine {:#x}", method_call.subchannel, method_call.argument); u32 engine = method_call.argument; @@ -109,12 +160,39 @@ void Puller::ProcessSemaphoreTriggerMethod(DmaPusher& dma_pusher) { if (op == GpuSemaphoreOperation::WriteLong) { const GPUVAddr sequence_address{regs.semaphore_address.SemaphoreAddress()}; const u32 payload = regs.semaphore_sequence; - dma_pusher.rasterizer->Query(sequence_address, VideoCommon::QueryType::Payload, VideoCommon::QueryPropertiesFlags::HasTimeout, payload, 0); - const u64 source_page = g_nx2_semaphore_source_page; - if (source_page != 0 && (sequence_address & ~GPUVAddr{0xFFF}) == source_page) { - const GPUVAddr mirror = g_nx2_semaphore_mirror_page + (sequence_address & 0xFFF); + // RELEASE_SIZE (bit 24) set means a 4-byte release without the timestamp; writing 16 + // bytes there clobbers the words next to it (Switch 2 packs semaphores tightly). + const auto properties = g_nx2_ampere && (regs.semaphore_trigger & (1U << 24)) != 0 + ? VideoCommon::QueryPropertiesFlags::IsAFence + : VideoCommon::QueryPropertiesFlags::HasTimeout; + dma_pusher.rasterizer->Query(sequence_address, VideoCommon::QueryType::Payload, properties, payload, 0); + if (const GPUVAddr mirror = g_nx2_ampere ? Nx2SemaphoreMirror(sequence_address) : 0; + mirror != 0) { LOG_DEBUG(HW_GPU, "Switch 2 semaphore {:#x}={:#x} mirrored to {:#x}", sequence_address, payload, mirror); - dma_pusher.rasterizer->Query(mirror, VideoCommon::QueryType::Payload, VideoCommon::QueryPropertiesFlags::HasTimeout, payload, 0); + dma_pusher.rasterizer->Query(mirror, VideoCommon::QueryType::Payload, properties, payload, 0); + } + } else if (g_nx2_ampere) { + // Switch 2 waits on semaphores the CPU (or another channel) releases, e.g. once it has + // written the commands this channel runs next: suspend the channel until the value is + // met instead of reading ahead (approach from changes.patch). + PendingAcquire acquire{ + .address = regs.semaphore_address.SemaphoreAddress(), + .value = regs.semaphore_sequence, + .operation = static_cast(op), + .since = std::chrono::steady_clock::now(), + .reported = false, + }; + if (op != GpuSemaphoreOperation::AcquireEqual && op != GpuSemaphoreOperation::AcquireGequal && + op != GpuSemaphoreOperation::AcquireMask) { + LOG_ERROR(HW_GPU, "Invalid semaphore operation {}", static_cast(op)); + return; + } + if (acquire.Matches(ReadSemaphoreWord(dma_pusher, acquire.address))) { + return; + } + dma_pusher.rasterizer->ReleaseFences(); + if (!acquire.Matches(ReadSemaphoreWord(dma_pusher, acquire.address))) { + pending_acquire = acquire; } } else { do { @@ -167,6 +245,36 @@ void Puller::ProcessSemaphoreAcquire(DmaPusher& dma_pusher) { } } +// Ampere SEM_EXECUTE (method 0x1B): operation in bits 2:0 on the semaphore at +// SEM_ADDR_HI:SEM_ADDR_LO (methods 0x18/0x17) with payload SEM_PAYLOAD_LO (0x19). +void Puller::ProcessSemaphoreExecute(DmaPusher& dma_pusher, u32 argument) { + const GPUVAddr address = (GPUVAddr{regs.reg_array[0x18] & 0x1FFFFFF} << 32) | + (regs.reg_array[0x17] & ~3U); + const u32 payload = regs.reg_array[0x19]; + const u32 operation = argument & 7; + if (operation == 1) { // RELEASE + dma_pusher.rasterizer->Query(address, VideoCommon::QueryType::Payload, + VideoCommon::QueryPropertiesFlags::IsAFence, payload, 0); + return; + } + // Acquires: check once after flushing pending fences; never block the GPU thread on them. + dma_pusher.rasterizer->ReleaseFences(); + const u32 word = dma_pusher.memory_manager.Read(address); + const bool ready = operation == 0 ? word == payload + : operation == 2 ? word >= payload + : operation == 3 ? static_cast(word - payload) >= 0 + : operation == 4 ? (word & payload) != 0 + : operation == 5 ? (word & payload) == 0 + : true; + if (!ready) { + static std::atomic reported{}; + if (reported.fetch_add(1) < 8) { + LOG_WARNING(HW_GPU, "SEM_EXECUTE op {} at {:#x}: {:#x} vs {:#x} not met, continuing", + operation, address, word, payload); + } + } +} + /// Calls a GPU puller method. void Puller::CallPullerMethod(DmaPusher& dma_pusher, const MethodCall& method_call) { static const bool nx2_trace = std::getenv("EDEN_NX2_PULLER_TRACE") != nullptr; @@ -221,11 +329,18 @@ void Puller::CallPullerMethod(DmaPusher& dma_pusher, const MethodCall& method_ca break; } case BufferMethods::SemaphoreAcquire: { - ProcessSemaphoreAcquire(dma_pusher); + // On Ampere host classes 0x17..0x1A are SEM_ADDR_LO/HI and SEM_PAYLOAD_LO/HI. + if (!g_nx2_ampere) { + ProcessSemaphoreAcquire(dma_pusher); + } break; } case BufferMethods::SemaphoreRelease: { - ProcessSemaphoreRelease(dma_pusher); + if (g_nx2_ampere) { + ProcessSemaphoreExecute(dma_pusher, method_call.argument); + } else { + ProcessSemaphoreRelease(dma_pusher); + } break; } case BufferMethods::Yield: { @@ -235,6 +350,10 @@ void Puller::CallPullerMethod(DmaPusher& dma_pusher, const MethodCall& method_ca } default: LOG_ERROR(HW_GPU, "Special puller engine method {:X} not implemented", method); + if (static bool dumped = false; g_nx2_ampere && !dumped) { + dumped = true; + dma_pusher.DumpCurrentEntry("unknown host method"); + } break; } } diff --git a/src/video_core/engines/puller.h b/src/video_core/engines/puller.h index e8994f5640..2b9e2c5df3 100644 --- a/src/video_core/engines/puller.h +++ b/src/video_core/engines/puller.h @@ -7,7 +7,9 @@ #pragma once #include +#include #include +#include #include #include "common/bit_field.h" #include "common/common_funcs.h" @@ -70,6 +72,24 @@ public: BitField<8, 24, u32> syncpoint_id; }; + /// Switch 2: a semaphore acquire this channel is suspended on. + struct PendingAcquire { + GPUVAddr address; + u32 value; + u32 operation; + std::chrono::steady_clock::time_point since; + bool reported; + + [[nodiscard]] bool Matches(u32 current) const; + }; + + [[nodiscard]] bool HasPendingAcquire() const { + return pending_acquire.has_value(); + } + [[nodiscard]] bool CanResumeAcquire(DmaPusher& dma_pusher); + /// Clears the pending acquire once it is met; false while the channel has to keep waiting. + bool ResumeAcquire(DmaPusher& dma_pusher); + void CallMethod(DmaPusher& dma_pusher, const MethodCall& method_call); void CallMultiMethod(DmaPusher& dma_pusher, u32 method, u32 subchannel, const u32* base_start, u32 amount, u32 methods_pending); void BindRasterizer(DmaPusher& dma_pusher, VideoCore::RasterizerInterface* rasterizer); @@ -125,8 +145,11 @@ private: void ProcessSemaphoreAcquire(DmaPusher& dma_pusher); void ProcessSemaphoreRelease(DmaPusher& dma_pusher); void ProcessSemaphoreTriggerMethod(DmaPusher& dma_pusher); + void ProcessSemaphoreExecute(DmaPusher& dma_pusher, u32 argument); [[nodiscard]] bool ExecuteMethodOnEngine(DmaPusher& dma_pusher, u32 method); + std::optional pending_acquire; + /// Mapping of command subchannels to their bound engine ids std::array bound_engines{}; diff --git a/src/video_core/gpu_thread.cpp b/src/video_core/gpu_thread.cpp index a01375e7be..97e18df499 100644 --- a/src/video_core/gpu_thread.cpp +++ b/src/video_core/gpu_thread.cpp @@ -4,6 +4,9 @@ // SPDX-FileCopyrightText: Copyright 2019 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include +#include + #include "common/assert.h" #include "common/scope_exit.h" #include "common/settings.h" @@ -36,7 +39,16 @@ void ThreadManager::StartThread(VideoCore::RendererBase& renderer, Core::Fronten auto current_context = context.Acquire(); CommandDataContainer next; while (!stop_token.stop_requested()) { - state.queue.PopWait(next, stop_token); + if (scheduler.HasPendingChannels()) { + // Suspended channels wait on semaphores the CPU releases: poll them meanwhile. + scheduler.ResumePending(system.GPU()); + if (!state.queue.TryPop(next)) { + std::this_thread::sleep_for(std::chrono::milliseconds{1}); + continue; + } + } else { + state.queue.PopWait(next, stop_token); + } if (stop_token.stop_requested()) { break; } diff --git a/src/video_core/memory_manager.cpp b/src/video_core/memory_manager.cpp index 0290500e7c..52e9db189d 100644 --- a/src/video_core/memory_manager.cpp +++ b/src/video_core/memory_manager.cpp @@ -251,7 +251,10 @@ T MemoryManager::Read(GPUVAddr addr) const { return value; } - ASSERT(false); + static std::atomic reported{}; + if (reported.fetch_add(1) < 16) { + LOG_ERROR(HW_GPU, "Read{} of unmapped GPU address {:#x}", sizeof(T) * 8, addr); + } return {}; } diff --git a/src/video_core/nx2.h b/src/video_core/nx2.h index 9fdf0a7c68..ef6619df3f 100644 --- a/src/video_core/nx2.h +++ b/src/video_core/nx2.h @@ -3,6 +3,7 @@ #pragma once +#include #include #include "common/common_types.h" @@ -13,10 +14,34 @@ namespace Tegra { // Ampere classes and the Turing+ macro ISA, which Eden's Maxwell engines don't implement. inline std::atomic g_nx2_ampere{false}; -// NVN2 queues release their progress values into one GPU page and wait (GPU acquires and CPU -// polls) on the same offsets in the page below it. Something on the console copies one into -// the other; until that is known, the puller mirrors releases that land in the source page. -inline std::atomic g_nx2_semaphore_source_page{0}; // GPU VA of the released-to page -inline std::atomic g_nx2_semaphore_mirror_page{0}; // GPU VA of the waited-on page +// NVN2 queues release their progress values into the second page of a two-page GPU-accessible +// region and GPU acquires read the same offsets in the first page (the CPU polls the second). +// Something on the console copies one into the other; until that is known, the puller mirrors +// releases that land in a second page. Every queue has such a region, so keep a small table. +inline constexpr size_t Nx2MaxSemaphoreRegions = 32; +inline std::array, Nx2MaxSemaphoreRegions> g_nx2_semaphore_regions{}; // starts + +inline void Nx2AddSemaphoreRegion(u64 start) { + for (auto& region : g_nx2_semaphore_regions) { + u64 expected = 0; + if (region.load() == start || region.compare_exchange_strong(expected, start)) { + return; + } + } +} + +/// GPU VA in the first page that mirrors a release at `address`, or 0. +inline u64 Nx2SemaphoreMirror(u64 address) { + for (const auto& region : g_nx2_semaphore_regions) { + const u64 start = region.load(std::memory_order_relaxed); + if (start == 0) { + break; + } + if (address >= start + 0x1000 && address < start + 0x2000) { + return address - 0x1000; + } + } + return 0; +} } // namespace Tegra diff --git a/src/video_core/query_cache/query_cache.h b/src/video_core/query_cache/query_cache.h index b8c1bc4f45..11b2f631a5 100644 --- a/src/video_core/query_cache/query_cache.h +++ b/src/video_core/query_cache/query_cache.h @@ -7,6 +7,7 @@ #pragma once #include +#include #include #include #include @@ -233,6 +234,10 @@ void QueryCacheBase::CounterReport(GPUVAddr addr, QueryType counter_type QueryPropertiesFlags flags, u32 payload, u32 subreport) { const bool has_timestamp = True(flags & QueryPropertiesFlags::HasTimeout); const bool is_fence = True(flags & QueryPropertiesFlags::IsAFence); + if (static const bool trace = std::getenv("EDEN_NX2_QUERY_TRACE") != nullptr; trace) { + LOG_INFO(HW_GPU, "nx2 report addr={:#x} type={} flags={:#x} payload={:#x}", addr, + static_cast(counter_type), static_cast(flags), payload); + } size_t streamer_id = static_cast(counter_type); auto* streamer = impl->streamers[streamer_id]; if (streamer == nullptr) [[unlikely]] { -- cgit v1.2.3 From 51207c73c889dbdc4a586a5da23314f6d33e4eb9 Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Tue, 27 Oct 2026 22:26:16 -0400 Subject: Switch 2: Turing macros, native compute, Ampere shader lowering, native tiling From changes.patch: the TU104/GA10x MME interpreter (shared macro context for the process's channels, also on the compute class), native compute with QMD 3.0 dependency scheduling, the Ampere-to-Maxwell shader lowering, native vertex stream sizes, and the console's GOB swizzle for textures, DMA copies and the host unswizzle shaders. Switch 2 channels turn all of it on at init. The lowering is widened for title 4 (all 32 of its programs lower): general TEX (any dimension/LOD/mask, operand split 1+1/2+1/2+2/4+n, any but consistent descriptor buffer), cube/array, SULD channel masks, 128-bit STG, indexed IPA, VOTE, SHFL with any immediate or predicate output, R2P/P2R byte select, uniform ULOP3/ULEA/UBREV/UFLO.SH/UISETP signed/UMOV URZ, IMNMX cbuf, IADD3 and FFMA uniform operands, FMUL /4 /8 x8, and BSSY/BSYNC/BREAK/WARPSYNC as no-ops. The capture stops at unmapped memory. EDEN_NX2_SHADER_DUMP dumps every native program; tools in C:\dev\nx2work test the lowering offline. --- src/shader_recompiler/CMakeLists.txt | 3 + src/shader_recompiler/environment.h | 2 + .../frontend/ampere/lower_to_maxwell.h | 1596 ++++++++++++++++++++ .../frontend/ampere/program_address_map.h | 29 + .../frontend/ampere/program_capture.h | 43 + .../frontend/ampere/program_override.h | 85 ++ .../frontend/maxwell/control_flow.cpp | 2 +- src/video_core/CMakeLists.txt | 3 + src/video_core/buffer_cache/buffer_cache.h | 8 +- src/video_core/control/channel_state.cpp | 23 + src/video_core/dirty_flags.cpp | 1 + src/video_core/engines/engine_upload.cpp | 20 + src/video_core/engines/engine_upload.h | 3 + src/video_core/engines/fermi_2d.cpp | 9 + src/video_core/engines/kepler_compute.cpp | 218 ++- src/video_core/engines/kepler_compute.h | 33 + src/video_core/engines/maxwell_3d.cpp | 31 +- src/video_core/engines/maxwell_3d.h | 18 + src/video_core/engines/maxwell_dma.cpp | 17 +- .../host_shaders/block_linear_unswizzle_2d.comp | 8 +- .../host_shaders/block_linear_unswizzle_3d.comp | 8 +- .../block_linear_unswizzle_3d_bcn.comp | 10 +- src/video_core/macro.cpp | 30 + src/video_core/macro.h | 19 + src/video_core/macro_turing.h | 213 +++ src/video_core/memory_manager.cpp | 4 + src/video_core/memory_manager.h | 10 + src/video_core/native_qmd.h | 112 ++ src/video_core/native_query_trace.h | 42 + .../renderer_vulkan/vk_texture_cache.cpp | 5 + src/video_core/shader_cache.cpp | 30 +- src/video_core/shader_environment.cpp | 184 ++- src/video_core/shader_environment.h | 11 + .../texture_cache/accelerated_swizzle.cpp | 2 + src/video_core/texture_cache/accelerated_swizzle.h | 2 + .../texture_cache/format_lookup_table.cpp | 1 + src/video_core/texture_cache/image_info.h | 12 + src/video_core/texture_cache/pending_download.h | 58 + src/video_core/texture_cache/texture_cache.h | 81 +- src/video_core/texture_cache/texture_cache_base.h | 7 +- src/video_core/texture_cache/util.cpp | 6 +- 41 files changed, 2930 insertions(+), 69 deletions(-) create mode 100644 src/shader_recompiler/frontend/ampere/lower_to_maxwell.h create mode 100644 src/shader_recompiler/frontend/ampere/program_address_map.h create mode 100644 src/shader_recompiler/frontend/ampere/program_capture.h create mode 100644 src/shader_recompiler/frontend/ampere/program_override.h create mode 100644 src/video_core/macro_turing.h create mode 100644 src/video_core/native_qmd.h create mode 100644 src/video_core/native_query_trace.h create mode 100644 src/video_core/texture_cache/pending_download.h (limited to 'src/video_core/memory_manager.cpp') diff --git a/src/shader_recompiler/CMakeLists.txt b/src/shader_recompiler/CMakeLists.txt index cc9ce95d92..39484a1bd5 100644 --- a/src/shader_recompiler/CMakeLists.txt +++ b/src/shader_recompiler/CMakeLists.txt @@ -58,6 +58,9 @@ add_library(shader_recompiler STATIC frontend/ir/type.h frontend/ir/value.cpp frontend/ir/value.h + frontend/ampere/lower_to_maxwell.h + frontend/ampere/program_override.h + frontend/ampere/program_address_map.h frontend/maxwell/control_flow.cpp frontend/maxwell/control_flow.h frontend/maxwell/decode.cpp diff --git a/src/shader_recompiler/environment.h b/src/shader_recompiler/environment.h index 184b56ebf3..67486430c8 100644 --- a/src/shader_recompiler/environment.h +++ b/src/shader_recompiler/environment.h @@ -74,6 +74,8 @@ public: [[nodiscard]] virtual u64 ReadInstruction(u32 address) = 0; + [[nodiscard]] virtual u32 TranslateProgramAddress(u32 address) const { return address; } + [[nodiscard]] virtual u32 ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) = 0; [[nodiscard]] virtual TextureType ReadTextureType(u32 raw_handle) = 0; diff --git a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h new file mode 100644 index 0000000000..524688f036 --- /dev/null +++ b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -0,0 +1,1596 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later +// Encoding reference: Mesa NAK sm70_encode.rs (MIT), Copyright 2022 Collabora, Ltd. +#pragma once + +#include +#include +#include +#include +#include +#include +#include +#include "common/common_types.h" + +namespace Shader::Ampere { + +struct Instruction { + std::array words; + + u32 Bits(u32 offset, u32 count) const { + u64 bits = words[offset / 32]; + if (offset / 32 + 1 < words.size()) { + bits |= u64{words[offset / 32 + 1]} << 32; + } + return static_cast((bits >> (offset % 32)) & ((u64{1} << count) - 1)); + } +}; +static_assert(sizeof(Instruction) == 16); + +struct LoweredProgram { + std::vector instructions; + std::vector source_instruction_offsets; + u32 source_size{}; + u32 required_texture_buffer = ~0U; + u32 error_offset{}; + std::string_view error; + + explicit operator bool() const { return error.empty(); } +}; + +// Lower only operations whose semantics are representable by the existing +// Maxwell frontend. These are new instructions, not reinterpretations of the +// native bytes. Scheduling/control words are added by the environment adapter. +inline LoweredProgram LowerToMaxwell(std::span code) { + LoweredProgram result; + struct Relocation { u32 output_index; s64 target; }; + std::vector relocations; + std::vector barrier_tokens; + std::vector call_targets; + std::vector call_links; + bool main_exited = false; + bool uses_uniform_bank = false; + for (const auto& instruction : code) { + const u32 op = instruction.Bits(0, 12); + uses_uniform_bank |= op == 0x3c4 || op == 0x2bd || op == 0xc0c || op == 0x986 || op == 0x9a8 || op == 0xc12 || op == 0xc17 || op == 0xe24 || op == 0xc24 || op == 0xd00 || op == 0xd09 || op == 0x9c3 || op == 0xc45 || op == 0x2a4 || op == 0x891 || op == 0x890 || op == 0x290 || op == 0x28c || op == 0x88c || op == 0x287 || op == 0x887 || op == 0x958 || op == 0xab9 || op == 0xabb || op == 0x899 || op == 0xc20 || op == 0xe21 || op == 0xe23 || op == 0xc02 || + op == 0x886 || op == 0x3c2 || op == 0xc82 || op == 0x882 || op == 0x292 || op == 0x892 || op == 0x291 || op == 0xc10 || op == 0xc23 || op == 0x526 || op == 0x2be || (op == 0x81c && instruction.Bits(67, 1)); + } + const auto address_of = [](u32 index) { return 0x50U + index * 8 + ((index + 1) / 3) * 8; }; + const auto fail = [&](u32 index, std::string_view reason) { + result.instructions.clear(); + result.source_instruction_offsets.clear(); + result.error_offset = index * sizeof(Instruction); + result.error = reason; + return result; + }; + const auto finish = [&](u32 index) { + for (const auto& relocation : relocations) { + if (relocation.target > index) return fail(index, "native branch or call leaves captured program"); + const u32 target = result.source_instruction_offsets[static_cast(relocation.target)]; + const s64 relative = s64{address_of(target)} - address_of(relocation.output_index) - 8; + if (relative < -0x800000 || relative > 0x7fffff) return fail(index, "lowered branch distance"); + result.instructions[relocation.output_index] |= (static_cast(relative) & 0xffffff) << 20; + } + result.source_size = (index + 1) * sizeof(Instruction); + return result; + }; + if (code.empty() || code.size() > 65536) return fail(0, "invalid native program size"); + for (u32 index = 0; index < code.size(); ++index) { + auto insn = code[index]; + result.source_instruction_offsets.push_back(static_cast(result.instructions.size())); + u32 opcode = insn.Bits(0, 12); + std::optional active_link; + u32 active_entry = 0; + if (main_exited) { + for (size_t call = 0; call < call_targets.size(); ++call) { + if (call_targets[call] <= index && call_targets[call] >= active_entry) { + active_entry = call_targets[call]; + active_link = call_links[call]; + } + } + if (active_link && insn.Bits(16, 8) == *active_link) + return fail(index, "native callee overwrites link register"); + } + // Scratch values and emulated uniform state occupy reserved high GPRs. + // Reject native accesses that could alias that state. + for (u32 bit : {16U, 24U, 32U, 64U}) { + if (bit == 32 && insn.Bits(9, 3) != 1 && opcode != 0xb60 && opcode != 0xb66) continue; + const u32 reg = insn.Bits(bit, 8); + if (reg >= (uses_uniform_bank ? 216U : 246U) && reg < 255) return fail(index, "native program uses lowering scratch registers"); + } + const auto uniform_register = [&](u32 reg) -> u32 { + if (reg < 16) return 216 + reg; + if (reg == 63) return 255; + return ~0U; + }; + if (opcode == 0x2a4 || opcode == 0x891 || opcode == 0x890 || opcode == 0x290 || + opcode == 0x292 || opcode == 0x892 || opcode == 0x291) { + if (insn.Bits(12, 4) != 7) return fail(index, "native uniform arithmetic predication"); + const u32 dst = uniform_register(insn.Bits(16, 8)); + const u32 a = uniform_register(insn.Bits(24, 8)); + const u32 c = uniform_register(insn.Bits(64, 8)); + if (dst == ~0U || a == ~0U || c == ~0U) return fail(index, "native uniform arithmetic register"); + insn.words[0] = (insn.words[0] & 0xffff) | (dst << 16) | (a << 24); + insn.words[2] = (insn.words[2] & ~0x080000ffU) | c; + if (opcode == 0x2a4 || opcode == 0x290 || opcode == 0x292 || opcode == 0x291) { + const u32 b = uniform_register(insn.Bits(32, 8)); + if (b == ~0U) return fail(index, "native uniform arithmetic second register"); + insn.words[1] = (insn.words[1] & ~0xffU) | b; + } + if ((opcode == 0x292 || opcode == 0x892) && insn.Bits(80, 4) != 0xe) + return fail(index, "native uniform logic predicate output"); + opcode = opcode == 0x2a4 ? 0x224 : opcode == 0x891 ? 0x811 : opcode == 0x890 ? 0x810 + : opcode == 0x292 ? 0x212 : opcode == 0x892 ? 0x812 : opcode == 0x291 ? 0x211 : 0x210; + } + const bool uniform_predication = opcode == 0xabb || opcode == 0xab9 || opcode == 0xc82 || opcode == 0x882; + const u32 native_predicate = insn.Bits(12, 4); + const bool guarded_uniform = uniform_predication && native_predicate != 7; + if (guarded_uniform && (native_predicate & 7) != 0) + return fail(index, "native uniform execution predicate"); + u64 predicate = u64{guarded_uniform ? ((native_predicate & 8) | 6) : native_predicate} << 16; + const u64 dest = insn.Bits(16, 8); + const auto base = [&](u64 opcode_bits) { return opcode_bits | predicate | dest; }; + const auto cbuf = [&]() -> u64 { + return (u64{insn.Bits(38, 16) / 4} << 20) | (u64{insn.Bits(54, 5)} << 34); + }; + const auto valid_cbuf = [&] { + return !insn.Bits(91, 1) && insn.Bits(54, 5) < 18 && + (insn.Bits(38, 16) % 4) == 0; + }; + std::string_view operand_error; + const auto emit = [&](u64 word) { result.instructions.push_back(word); }; + const auto save_predicate = [&] { + emit(0x38e8000000070000ULL | 253 | (u64{255} << 8) | (u64{0x40} << 20)); + }; + const auto restore_predicate = [&] { + emit(0x38f0000000070000ULL | (u64{253} << 8) | (u64{0x40} << 20)); + }; + if (guarded_uniform) { + save_predicate(); + emit(0x38f0000000070000ULL | (u64{239} << 8) | (u64{0x40} << 20)); + } + const auto mov_imm = [&](u32 reg, u32 value) { + emit(0x0100000000000000ULL | predicate | reg | (u64{15} << 12) | (u64{value} << 20)); + }; + const auto materialize = [&](u32 kind, u32 reg, u32 scratch, bool absolute, bool negate) -> u32 { + if (kind == 3) { + if (reg < 16) reg += 216; + else if (reg == 63) reg = 255; + else { operand_error = "native uniform source register"; return 255; } + } else if (kind == 1) { + mov_imm(scratch, insn.words[1]); + reg = scratch; + } else if (kind == 2) { + if (!valid_cbuf()) { operand_error = "native constant buffer address"; return 255; } + emit(0x4c98000000000000ULL | predicate | scratch | (u64{15} << 39) | cbuf()); + reg = scratch; + } + if (absolute || negate) { + if (reg != scratch) emit(0x5c98000000000000ULL | predicate | scratch | (u64{reg} << 20) | (u64{15} << 39)); + if (absolute) emit(0x0400000000000000ULL | predicate | scratch | (u64{scratch} << 8) | (u64{0x7fffffff} << 20)); + if (negate) emit(0x0400000000000000ULL | predicate | scratch | (u64{scratch} << 8) | (u64{0x80000000} << 20) | (u64{2} << 53)); + reg = scratch; + } + return reg; + }; + const auto operand = [&](u32 which, bool floating = true) -> u32 { + const u32 form = insn.Bits(9, 3); + u32 reg{}, kind{}, abs{}, neg{}; + if (which == 0) { + reg = insn.Bits(24, 8); abs = insn.Bits(73, 1); neg = insn.Bits(72, 1); + } else { + if (form != 1 && form != 2 && form != 3 && form != 4 && form != 5 && form != 6 && form != 7) { + operand_error = "native ALU source form"; return 255; + } + const bool wide = which == 1 ? (form == 4 || form == 5 || form == 6) : (form == 2 || form == 3 || form == 7); + if (wide) { + kind = form == 6 || form == 7 ? 3 : (form == 4 || form == 2) ? 1 : 2; + if (kind == 3) reg = insn.Bits(32, 8); + if (kind == 2 || kind == 3) { abs = insn.Bits(62, 1); neg = insn.Bits(63, 1); } + } else { + const bool upper = which == 1 ? (form == 2 || form == 3 || form == 7) : true; + reg = insn.Bits(upper ? 64 : 32, 8); + abs = insn.Bits(upper ? 74 : 62, 1); neg = insn.Bits(upper ? 75 : 63, 1); + } + } + return materialize(kind, reg, 250 + which, floating && abs, floating && neg); + }; + struct HalfSource { u32 reg, swizzle; bool absolute, negate; }; + const auto half_source = [&](u32 which) -> HalfSource { + if (which == 0) return {insn.Bits(24, 8), insn.Bits(74, 2), + insn.Bits(73, 1) != 0, insn.Bits(72, 1) != 0}; + const u32 form = insn.Bits(9, 3); + const bool wide = which == 1 ? (form == 4 || form == 5) : (form == 2 || form == 3); + if (wide) { + const bool immediate = form == 2 || form == 4; + const u32 reg = materialize(immediate ? 1 : 2, 255, 250 + which, false, false); + return {reg, immediate ? 0U : insn.Bits(60, 2), + !immediate && insn.Bits(62, 1), !immediate && insn.Bits(63, 1)}; + } + const bool upper = which == 2 || form == 2 || form == 3; + return {insn.Bits(upper ? 64 : 32, 8), insn.Bits(upper ? 81 : 60, 2), + insn.Bits(upper ? 83 : 62, 1) != 0, insn.Bits(upper ? 84 : 63, 1) != 0}; + }; + const auto extract_half = [&](HalfSource source, u32 lane, u32 target) { + const bool high = source.swizzle == 3 || (source.swizzle == 0 && lane == 1); + emit(0x5ca8000000000000ULL | predicate | target | (u64{2} << 8) | + (u64{source.swizzle == 1 ? 2U : 1U} << 10) | (u64{source.reg} << 20) | + (u64{high} << 41) | (u64{source.negate} << 45) | (u64{source.absolute} << 49)); + }; + const auto xmad = [&](u32 target, u32 a, u32 b, u32 c, bool a_high, + bool b_high, u32 mode, bool immediate) { + emit((immediate ? 0x3600000000000000ULL : 0x5b00000000000000ULL) | + predicate | target | (u64{a} << 8) | + (u64{immediate ? ((b >> (b_high ? 16 : 0)) & 0xffff) : b} << 20) | + (u64{!immediate && b_high} << 35) | (u64{c} << 39) | + (u64{mode} << 50) | (u64{a_high} << 53)); + }; + const auto multiply_low = [&](u32 a, u32 b, u32 c, bool immediate) { + xmad(250, a, b, c, false, false, 0, immediate); + xmad(251, a, b, 255, true, false, 0, immediate); + xmad(251, a, b, 251, false, true, 0, immediate); + emit(base(0x5c18000000000000ULL) | (u64{251} << 8) | + (u64{250} << 20) | (u64{16} << 39)); + }; + switch (opcode) { + case 0x304: { // F2F.F16.F32, scalar round-to-nearest conversion. + if (insn.words[2] != 0x00200800 || insn.Bits(24, 8) || + (insn.words[1] & ~0xc00000ffU)) + return fail(index, "native scalar float conversion mode"); + emit(base(0x5ca8000000000000ULL) | (u64{1} << 8) | (u64{2} << 10) | + (u64{insn.Bits(32, 8)} << 20) | (u64{insn.Bits(62, 1)} << 49) | + (u64{insn.Bits(63, 1)} << 45)); + break; + } + case 0x23e: case 0xa3e: { // F2FP.PACK_AB / MERGE_C, round-to-nearest + if (insn.Bits(62, 2) || insn.Bits(73, 5) || insn.Bits(79, 2) || + insn.Bits(81, 10) || insn.Bits(91, 1)) { + return fail(index, "native float packing modifiers"); + } + const bool merge = insn.Bits(78, 1); + if (insn.Bits(72, 1)) return fail(index, "native float packing selector"); + const u32 a = merge ? operand(2, false) : operand(0, false); + const u32 b = operand(1, false); + // Canonical F2F writes the converted half in its low 16 bits. + if (merge) { + emit(0x5c98000000000000ULL | predicate | 250 | (u64{a} << 20) | (u64{15} << 39)); + } else { + emit(0x5ca8000000000000ULL | predicate | 250 | (u64{1} << 8) | + (u64{2} << 10) | (u64{a} << 20)); + } + emit(0x5ca8000000000000ULL | predicate | 251 | (u64{1} << 8) | + (u64{2} << 10) | (u64{b} << 20)); + emit(base(0x5c18000000000000ULL) | (u64{250} << 8) | (u64{251} << 20) | (u64{16} << 39)); + break; + } + case 0x232: case 0x832: case 0xa32: { // HMUL2 + if (insn.Bits(78, 2) || insn.Bits(81, 10) || insn.Bits(91, 1)) + return fail(index, "native half multiply mode"); + const auto a = half_source(0), b = half_source(1); + const u32 precision = insn.Bits(76, 1) ? 2 : insn.Bits(80, 1); + emit(base(0x5d08000000000000ULL) | (u64{a.reg} << 8) | (u64{b.reg} << 20) | + (u64{b.swizzle} << 28) | (u64{b.absolute} << 30) | (u64{a.negate != b.negate} << 31) | + (u64{insn.Bits(77, 1)} << 32) | (u64{precision} << 39) | + (u64{a.absolute} << 44) | (u64{a.swizzle} << 47)); + break; + } + case 0x230: case 0x430: case 0x630: { // HADD2 + if (insn.Bits(76, 1) || insn.Bits(79, 1) || insn.Bits(81, 10) || insn.Bits(91, 1)) + return fail(index, "native half add mode"); + const auto a = half_source(0); + const auto b = half_source(opcode == 0x230 ? 1 : 2); + if (insn.Bits(78, 1)) { + if (a.reg != 255 || !a.negate || a.absolute || a.swizzle != 0 || + b.swizzle < 2 || insn.Bits(77, 1) || insn.Bits(80, 1)) + return fail(index, "native float-output half add mode"); + extract_half(b, 0, static_cast(dest)); + } else { + emit(base(0x5d10000000000000ULL) | (u64{a.reg} << 8) | (u64{b.reg} << 20) | + (u64{b.swizzle} << 28) | (u64{b.absolute} << 30) | (u64{b.negate} << 31) | + (u64{insn.Bits(77, 1)} << 32) | (u64{insn.Bits(80, 1)} << 39) | + (u64{a.negate} << 43) | (u64{a.absolute} << 44) | (u64{a.swizzle} << 47)); + } + break; + } + case 0x231: case 0x431: case 0x631: case 0x831: case 0xa31: { // HFMA2 + const bool relu = insn.Bits(79, 1) != 0; + if (insn.Bits(78, 1) || insn.Bits(85, 1) || insn.Bits(91, 1) || + (relu ? (insn.Bits(87, 4) != 7 || + insn.Bits(77, 1)) : insn.Bits(87, 4) != 0)) + return fail(index, "native half fused multiply-add mode"); + if (relu && native_predicate != 7) { + // Guard the whole expansion before its temporary P6 updates. + relocations.push_back({static_cast(result.instructions.size()), index + 1}); + emit(0xe24000000000000fULL | (u64{native_predicate ^ 8U} << 16)); + predicate = u64{7} << 16; + } + auto a = half_source(0); + auto b = half_source(1); + const auto c = half_source(2); + if (insn.Bits(86, 1)) { // H0_NH1: negate only the upper packed component. + if (b.swizzle != 0 || b.absolute || b.reg >= 250) + return fail(index, "native half mixed-sign source mode"); + emit(0x0400000000000000ULL | predicate | 251 | (u64{b.reg} << 8) | + (u64{0x80000000} << 20) | (u64{2} << 53)); + b.reg = 251; + } + if (b.absolute || c.absolute || (a.absolute && a.reg >= 248)) + return fail(index, "native half fused multiply-add absolute source"); + if (a.absolute) { + emit(0x0400000000000000ULL | predicate | 250 | (u64{a.reg} << 8) | + (u64{a.swizzle == 1 ? 0x7fffffffU : 0x7fff7fffU} << 20)); + a.reg = 250; + a.absolute = false; + } + const u32 precision = insn.Bits(76, 1) ? 2 : insn.Bits(80, 1); + emit(0x5d00000000000000ULL | predicate | (relu ? 254U : dest) | (u64{a.reg} << 8) | (u64{b.reg} << 20) | + (u64{b.swizzle} << 28) | (u64{c.negate} << 30) | (u64{a.negate != b.negate} << 31) | + (u64{insn.Bits(77, 1)} << 32) | (u64{c.swizzle} << 35) | + (u64{precision} << 37) | (u64{c.reg} << 39) | (u64{a.swizzle} << 47)); + if (relu) { + // Clamp the rounded binary16 lanes, preserving positive infinity + // and canonicalizing either sign of NaN without saturating to one. + save_predicate(); + for (u32 lane = 0; lane < 2; ++lane) { + emit(0x3828000000070000ULL | 251 | (u64{254} << 8) | (u64{lane * 16} << 20)); + emit(0x0400000000070000ULL | 251 | (u64{251} << 8) | (u64{0xffff} << 20)); + emit(0x0400000000070000ULL | 252 | (u64{251} << 8) | (u64{0x7fff} << 20)); + mov_imm(250, 0x7c00); + emit(0x5b60000000070000ULL | 7 | (u64{6} << 3) | (u64{250} << 8) | + (u64{252} << 20) | (u64{7} << 39) | (u64{1} << 49)); + emit(0x3828000000070000ULL | 252 | (u64{251} << 8) | (u64{15} << 20)); + mov_imm(250, 1); + emit(0x5c10000000070000ULL | 252 | (u64{252} << 8) | (u64{250} << 20) | (u64{1} << 48)); + emit(0x5c40000000070000ULL | 251 | (u64{251} << 8) | (u64{252} << 20) | (u64{7} << 48)); + mov_imm(250, 0x7fff); + emit(0x5ca0000000070000ULL | 251 | (u64{250} << 8) | (u64{251} << 20) | (u64{6} << 39)); + if (lane == 0) { + emit(base(0x5c98000000000000ULL) | (u64{251} << 20) | (u64{15} << 39)); + } else { + emit(0x3848000000070000ULL | 251 | (u64{251} << 8) | (u64{16} << 20)); + emit(base(0x5c40000000000000ULL) | (dest << 8) | (u64{251} << 20) | + (u64{1} << 41) | (u64{7} << 48)); + } + } + restore_predicate(); + } + break; + } + case 0x240: case 0x840: { // HMNMX2, exact half promotion for min/max + if (insn.Bits(76, 11) || insn.Bits(91, 1)) + return fail(index, "native half min/max mode"); + const auto a = half_source(0); + const auto b = opcode == 0x840 ? HalfSource{255, 0, false, false} : half_source(1); + if (a.swizzle == 1 || b.swizzle == 1 || b.reg >= 250 && b.reg != 255) + return fail(index, "native half min/max source mode"); + for (u32 lane = 0; lane < 2; ++lane) { + extract_half(a, lane, 250); + if (opcode == 0x840) { + const u32 half = (insn.words[1] >> (lane * 16)) & 0xffff; + mov_imm(251, half); + extract_half(HalfSource{251, 2, false, false}, 0, 251); + } else { + extract_half(b, lane, 251); + } + const u32 target = lane == 0 ? 252 : 250; + emit(0x5c60000000000000ULL | predicate | target | (u64{250} << 8) | (u64{251} << 20) | + (u64{insn.Bits(87, 3)} << 39) | (u64{insn.Bits(90, 1)} << 42)); + emit(0x5ca8000000000000ULL | predicate | target | (u64{1} << 8) | + (u64{2} << 10) | (u64{target} << 20)); + } + emit(base(0x5c18000000000000ULL) | (u64{250} << 8) | (u64{252} << 20) | (u64{16} << 39)); + break; + } + case 0xc24: case 0xe24: { // IMAD low word with a uniform multiplier or addend. + if ((insn.words[2] & ~0x080002ffU) != 0x078e0000 || (insn.words[1] & ~0xffU)) + return fail(index, "native uniform-addend multiply mode"); + const u32 a = operand(0, false), b = operand(1, false), c = operand(2, false); + multiply_low(a, b, c, false); + break; + } + case 0x224: { // IMAD register / zero-product move + if (insn.Bits(24, 8) == 255 && insn.words[1] == 255 && insn.Bits(75, 1) && + (insn.words[2] & ~0xaffU) == 0x078e0000) { + emit(base(0x5c10000000000000ULL) | (u64{255} << 8) | + (u64{insn.Bits(64, 8)} << 20) | (u64{1} << 48)); + break; + } + if (insn.Bits(24, 8) == 255 && insn.words[1] == 255 && + (insn.words[2] & ~0xffU) == 0x078e0000) { + emit(base(0x5c98000000000000ULL) | (u64{15} << 39) | (u64{insn.Bits(64, 8)} << 20)); + break; + } + if ((insn.words[2] & ~0x2ffU) != 0x078e0000 || (insn.words[1] & ~0xffU)) + return fail(index, "native integer register multiply mode"); + multiply_low(insn.Bits(24, 8), insn.Bits(32, 8), insn.Bits(64, 8), false); + break; + } + case 0x827: { // IMAD.HI.U32, captured zero addend + if (insn.words[2] != 0x078e00ff) return fail(index, "native high multiply mode"); + const u32 a = insn.Bits(24, 8), b = insn.words[1]; + // Base-2^16 long multiplication. Each intermediate sum fits U32. + xmad(250, a, b, 255, false, false, 0, true); + xmad(251, a, b, 250, true, false, 2, true); + xmad(250, a, b, 251, false, true, 1, true); + xmad(252, a, b, 251, true, true, 2, true); + emit(0x3828000000000000ULL | predicate | 250 | (u64{250} << 8) | (u64{16} << 20)); + emit(base(0x5c10000000000000ULL) | (u64{250} << 8) | (u64{252} << 20)); + break; + } + case 0xa24: { // IMAD, constant-buffer multiplier + if ((insn.words[2] & ~0x2ffU) != 0x078e0000 || !valid_cbuf() || insn.Bits(62, 2)) + return fail(index, "native integer constant-buffer multiply mode"); + emit(0x4c98000000000000ULL | predicate | 252 | (u64{15} << 39) | cbuf()); + multiply_low(insn.Bits(24, 8), 252, insn.Bits(64, 8), false); + break; + } + case 0x424: case 0x624: { // IMAD.MOV.U32: zero product plus immediate/cbuf + if (opcode == 0x624 && insn.Bits(24, 8) == 255 && + (insn.words[2] & ~0x07800000U) == 0x000e06ffU) { + const u32 source = operand(2, false); + mov_imm(250, 1); + emit(0x5ca0000000000000ULL | predicate | 250 | (u64{250} << 8) | (u64{255} << 20) | + (u64{insn.Bits(87, 3)} << 39) | (u64{insn.Bits(90, 1)} << 42)); + emit(base(0x5c10000000000000ULL) | (u64{source} << 8) | (u64{250} << 20)); + break; + } + if (insn.Bits(24, 8) != 255 || insn.Bits(64, 8) != 255 || insn.words[2] != 0x078e00ff) { + if ((insn.words[2] & ~0x2ffU) != 0x078e0000 || (opcode == 0x624 && insn.Bits(62, 2))) + return fail(index, "native multiply wide addend mode"); + const u32 a = operand(0, false), b = operand(1, false), c = operand(2, false); + multiply_low(a, b, c, false); + break; + } + if (insn.Bits(24, 8) != 255 || insn.Bits(64, 8) != 255 || + insn.words[2] != 0x078e00ff) + return fail(index, "native integer move mode"); + if (opcode == 0x624) { + if (!valid_cbuf() || insn.Bits(62, 2)) + return fail(index, "native integer move constant buffer"); + emit(base(0x4c98000000000000ULL) | (u64{15} << 39) | cbuf()); + } else { + emit(base(0x0100000000000000ULL) | (u64{15} << 12) | + (u64{insn.words[1]} << 20)); + } + break; + } + case 0x824: { // Captured low-word IMAD forms + if (insn.words[1] == 1 && (insn.words[2] & 0xffffff00U) == 0x078e0200) { + emit(base(0x5c10000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{insn.Bits(64, 8)} << 20)); + break; + } + if (insn.words[1] != 48 || insn.words[2] != 0x078e02ff) { + if ((insn.words[2] & ~0x2ffU) != 0x078e0000) + return fail(index, "native integer multiply mode"); + multiply_low(insn.Bits(24, 8), insn.words[1], insn.Bits(64, 8), true); + break; + } + const u64 source = insn.Bits(24, 8); + emit(0x5c18000000000000ULL | predicate | 250 | (source << 8) | (source << 20) | (u64{1} << 39)); + emit(base(0x3848000000000000ULL) | (u64{250} << 8) | (u64{4} << 20)); + break; + } + case 0x806: { // VOTE: register ballot, ordinary predicate input and output. + if (insn.words[1] || insn.Bits(72, 2) > 2 || (insn.words[2] & ~0x078e0300U)) + return fail(index, "native vote mode"); + emit(base(0x50d8000000000000ULL) | (u64{insn.Bits(87, 4)} << 39) | + (u64{insn.Bits(81, 3)} << 45) | (u64{insn.Bits(72, 2)} << 48)); + break; + } + case 0x886: { // VOTEU ballot and independent uniform predicate output. + const u32 output = insn.Bits(81, 3), target = insn.Bits(16, 8); + if ((target >= 16 && target != 63) || insn.words[1] || + (output > 1 && output != 7) || insn.Bits(72, 2) > 2 || native_predicate != 7) + return fail(index, "native uniform vote mode"); + const u32 temporary = output == 7 ? 7 : 6 - output; + const u32 mask = 1U << temporary; + if (output != 7) emit(0x38e8000000070000ULL | 253 | (u64{255} << 8) | (u64{mask} << 20)); + emit(0x50d8000000070000ULL | uniform_register(target) | + (u64{insn.Bits(87, 4)} << 39) | (u64{temporary} << 45) | (u64{insn.Bits(72, 2)} << 48)); + if (output != 7) { + emit(0x38e8000000070000ULL | 239 | (u64{239} << 8) | (u64{mask} << 20)); + emit(0x38f0000000070000ULL | (u64{253} << 8) | (u64{mask} << 20)); + } + break; + } + case 0x3c2: { // R2UR: broadcast the first lane satisfying the execution predicate. + if (dest >= 16 || insn.words[1] || insn.words[2] != 0x000e0000) + return fail(index, "native register-to-uniform mode"); + save_predicate(); + emit(0x50d8000000070000ULL | 250 | (u64{native_predicate} << 39) | + (u64{6} << 45) | (u64{1} << 48)); + emit(0x5c10000000070000ULL | 251 | (u64{250} << 8) | (u64{255} << 20) | (u64{1} << 49)); + emit(0x5c40000000070000ULL | 250 | (u64{250} << 8) | (u64{251} << 20) | (u64{7} << 48)); + emit(0x5c30000000070000ULL | 250 | (u64{250} << 20)); + emit(0xef10000000070000ULL | 251 | (u64{insn.Bits(24, 8)} << 8) | + (u64{250} << 20) | (u64{1} << 29) | (u64{31} << 34) | (u64{7} << 48)); + emit(0x5ca0000000070000ULL | (dest + 216) | (u64{251} << 8) | + (u64{dest + 216} << 20) | (u64{6} << 39)); + restore_predicate(); + break; + } + case 0x81c: { // Observed predicate LUT constants and UP0-to-predicate copy. + const u32 target = insn.Bits(81, 3); + const u32 lut = insn.Bits(64, 3) | (insn.Bits(72, 5) << 3); + if ((native_predicate != 7 && insn.Bits(67, 1)) || target == 7 || insn.Bits(84, 3) != 7 || + insn.Bits(16, 8) || insn.words[1]) + return fail(index, "native predicate LUT mode"); + if (!insn.Bits(67, 1) && (lut == 0xfe || lut == 0xa8)) { + emit(0x5090000000000000ULL | predicate | (u64{target} << 3) | 7 | + (u64{insn.Bits(87, 4)} << 12) | (u64{1} << 24) | + (u64{insn.Bits(77, 4)} << 29) | (u64{insn.Bits(68, 4)} << 39) | + (u64{lut == 0xfe} << 45)); + break; + } + if (!insn.Bits(67, 1) && (lut == 0x80 || + (lut == 8 && insn.Bits(68, 4) == 7 && insn.Bits(77, 4) == 7))) { + emit(0x5090000000000000ULL | predicate | (u64{target} << 3) | 7 | + (u64{insn.Bits(87, 4) ^ (lut == 8 ? 8U : 0U)} << 12) | (u64{insn.Bits(77, 4)} << 29) | + (u64{insn.Bits(68, 4)} << 39)); + break; + } + if (insn.Bits(77, 4) != 7 || insn.Bits(87, 4) != 7) + return fail(index, "native predicate LUT source mode"); + if (insn.Bits(67, 1) && insn.Bits(68, 4) == 0 && lut == 0x80) { + emit(0x3828000000070000ULL | 250 | (u64{239} << 8) | (u64{6} << 20)); + emit(0x3848000000070000ULL | 250 | (u64{250} << 8) | (u64{target} << 20)); + } else if (!insn.Bits(67, 1) && insn.Bits(68, 4) == 7 && (lut == 0x80 || lut == 8)) { + emit(0x010000000007f000ULL | 250 | ((lut == 0x80 ? u64{1} << target : 0ULL) << 20)); + } else return fail(index, "native predicate LUT truth table"); + emit(0x38f0000000070000ULL | (u64{250} << 8) | (u64{1} << (target + 20))); + break; + } + case 0x882: { // UMOV immediate. + if (dest >= 16 || insn.words[2]) return fail(index, "native uniform immediate move mode"); + mov_imm(static_cast(dest) + 216, insn.words[1]); + break; + } + case 0xc82: { // UMOV UR, UR (or URZ), including uniform execution predicates. + const u32 source = insn.words[1] < 64 ? uniform_register(insn.words[1]) : ~0U; + if (dest >= 16 || source == ~0U || insn.words[2] != 0x08000000) + return fail(index, "native uniform move mode"); + emit(0x5c98000000000000ULL | predicate | (dest + 216) | + (u64{source} << 20) | (u64{15} << 39)); + break; + } + case 0xab9: case 0xabb: { // ULDC U32/U64, direct or indexed. + const u32 uniform_dest = insn.Bits(16, 8), index_reg = insn.Bits(24, 8); + const u32 width = insn.Bits(73, 3); + if (uniform_dest >= 16 || (width != 4 && width != 5) || + (width == 5 && ((uniform_dest & 1) || uniform_dest >= 15)) || + insn.Bits(54, 5) >= 18 || insn.Bits(38, 16) % 4 || insn.Bits(62, 2) || + (insn.words[2] & ~0xe00U) != (opcode == 0xab9 ? 0U : 0x08000000U) || + (opcode == 0xabb && index_reg >= 16)) + return fail(index, "native uniform constant load mode"); + if (opcode == 0xab9 && width == 4) { + emit(0x4c98000000000000ULL | predicate | (uniform_dest + 216) | (u64{15} << 39) | cbuf()); + } else { + emit(0xef90000000000000ULL | predicate | (uniform_dest + 216) | + (u64{opcode == 0xab9 ? 255U : index_reg + 216} << 8) | (u64{insn.Bits(38, 16)} << 20) | + (u64{insn.Bits(54, 5)} << 36) | (u64{width} << 48)); + } + break; + } + case 0x899: { // USHF.L.U32; accept a zero word outside the selected half. + const u32 dst = insn.Bits(16, 8), low = insn.Bits(24, 8), high = insn.Bits(64, 8); + if (insn.words[2] == (0x08011400U | high) && low == 63 && high < 16 && + dst < 16 && native_predicate == 7 && insn.words[1] <= 32) { + emit(0x3828000000000000ULL | predicate | (dst + 216) | + (u64{high + 216} << 8) | (u64{insn.words[1]} << 20) | (u64{1} << 48)); + break; + } + const bool high_result = insn.words[2] == (0x08010600U | high) && low == 63 && high < 16; + const bool low_result = insn.words[2] == 0x0800063f && low < 16; + if (dst >= 16 || native_predicate != 7 || insn.words[1] > 32 || (!high_result && !low_result)) + return fail(index, "native uniform shift mode"); + emit(0x3848000000000000ULL | predicate | (dst + 216) | + (u64{(high_result ? high : low) + 216} << 8) | (u64{insn.words[1]} << 20)); + break; + } + case 0xb82: { // LDC.64 indexed byte offset + const u32 width = insn.Bits(73, 3); + if ((width != 4 && width != 5) || insn.Bits(78, 2) || insn.Bits(91, 1) || + insn.Bits(54, 5) >= 18 || + (width == 5 && ((dest & 1) || dest >= 254))) return fail(index, "native indexed constant load mode"); + emit(base(0xef90000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{insn.Bits(38, 16)} << 20) | (u64{insn.Bits(54, 5)} << 36) | (u64{width} << 48)); + break; + } + case 0xf89: { // SHFL with immediate lane and segmentation/clamp mask, any mode. + const u32 lane = insn.Bits(53, 5), mask = insn.Bits(40, 13), mode = insn.Bits(58, 2); + if (insn.words[2] != 0x000e0000 || insn.Bits(32, 8) || insn.Bits(60, 4)) + return fail(index, "native immediate shuffle mode"); + emit(base(0xef10000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{lane} << 20) | (u64{1} << 28) | (u64{1} << 29) | + (u64{mode} << 30) | (u64{mask} << 34) | (u64{7} << 48)); + break; + } + case 0x822: { // FSWZADD, bounded quad derivative sign patterns. + const u32 swizzle = insn.words[1]; + if ((swizzle != 0x66 && swizzle != 0x5a) || insn.words[2] > 255) + return fail(index, "native swizzled add mode"); + // Maxwell encodes operand signs rather than the native operation enum. + emit(base(0x50f8000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{insn.Bits(64, 8)} << 20) | (u64{swizzle == 0x66 ? 0x99U : 0xa5U} << 28)); + break; + } + case 0xb6c: { // TXD ARRAY_2D, four channels and consecutive split destinations. + const u32 binding = insn.Bits(54, 5), offset = insn.Bits(40, 14); + const u32 coordinates = insn.Bits(24, 8), derivatives = insn.Bits(32, 8); + if ((insn.words[2] & ~0xffU) != 0x001e0f00 || insn.Bits(59, 2) || + insn.Bits(61, 3) != 5 || binding >= 18 || offset > 0x1fff || + (dest & 1) || dest >= 248 || insn.Bits(64, 8) != dest + 2 || + coordinates >= 248 || derivatives >= 247) + return fail(index, "native array gradient texture mode"); + if (result.required_texture_buffer != ~0U && result.required_texture_buffer != binding) + return fail(index, "native texture descriptor buffer conflict"); + result.required_texture_buffer = binding; + emit(base(0xde00000000000000ULL) | (u64{coordinates} << 8) | + (u64{derivatives} << 20) | (u64{3} << 28) | (u64{15} << 31) | + (u64{offset} << 36) | (u64{insn.Bits(90, 1)} << 49) | (u64{7} << 51)); + break; + } + case 0x589: { // SHFL.IDX, register lane and immediate segmentation mask. + if ((insn.words[2] & ~0x000e0000U) || insn.Bits(53, 11)) + return fail(index, "native shuffle mode"); + emit(0xef10000000000000ULL | predicate | dest | + (u64{insn.Bits(24, 8)} << 8) | (u64{insn.Bits(32, 8)} << 20) | + (u64{1} << 29) | (u64{insn.Bits(40, 13)} << 34) | (u64{insn.Bits(81, 3)} << 48)); + break; + } + case 0x98e: { // RED.E.ADD.STRONG.GPU U32, direct register address. + if (insn.words[2] != 0x0010e100 || insn.Bits(40, 24) || (insn.Bits(24, 8) & 1)) + return fail(index, "native global reduction mode"); + emit(0xebf8000000000000ULL | predicate | insn.Bits(32, 8) | + (u64{insn.Bits(24, 8)} << 8) | (u64{1} << 48)); + break; + } + case 0x2bd: // UFLO.U32, the highest set-bit index (.SH, bit 74: its shift amount). + if (dest >= 16 || insn.words[1] >= 16 || (insn.words[2] & ~0x400U) != 0x080e0000 || native_predicate != 7) + return fail(index, "native uniform leading bit mode"); + emit(0x5c30000000000000ULL | predicate | (dest + 216) | (u64{insn.words[1] + 216} << 20) | + (u64{insn.Bits(74, 1)} << 41)); + break; + case 0x2be: // UBREV, as a bit-reversing zero shift like BREV. + if (dest >= 16 || insn.words[1] >= 16 || insn.words[2] != 0x08000000 || native_predicate != 7) + return fail(index, "native uniform bit reverse mode"); + emit(0x3828000000000000ULL | predicate | (dest + 216) | (u64{insn.words[1] + 216} << 8) | + (u64{1} << 39) | (u64{1} << 40)); + break; + case 0x3a1: { // MATCH.ALL.U64 with only the all-equal predicate retained. + const u32 source = insn.Bits(24, 8), target = insn.Bits(81, 3); + if (dest != 255 || source >= 254 || (source & 1) || target >= 6 || + native_predicate != 7 || insn.words[1] || (insn.words[2] & ~0xe0000U) != 0x200) + return fail(index, "native warp match mode"); + save_predicate(); + emit(0x50d8000000070000ULL | 250 | (u64{7} << 39) | (u64{7} << 45) | (u64{1} << 48)); + emit(0x5c10000000070000ULL | 251 | (u64{250} << 8) | (u64{255} << 20) | (u64{1} << 49)); + emit(0x5c40000000070000ULL | 250 | (u64{250} << 8) | (u64{251} << 20) | (u64{7} << 48)); + emit(0x5c30000000070000ULL | 250 | (u64{250} << 20)); + for (u32 word = 0; word < 2; ++word) + emit(0xef10000000070000ULL | (251 + word) | (u64{source + word} << 8) | + (u64{250} << 20) | (u64{1} << 29) | (u64{31} << 34) | (u64{7} << 48)); + emit(0x5b60000000070000ULL | 7 | (u64{6} << 3) | (u64{source} << 8) | + (u64{251} << 20) | (u64{7} << 39) | (u64{2} << 49)); + emit(0x5b60000000070000ULL | 7 | (u64{6} << 3) | (u64{source + 1} << 8) | + (u64{252} << 20) | (u64{6} << 39) | (u64{2} << 49)); + emit(0x50d8000000070000ULL | 255 | (u64{6} << 39) | (u64{target} << 45)); + restore_predicate(); + break; + } + case 0x3c4: { // REDUX.SUM over the actual active lanes, modulo 2^32. + const u32 source = insn.Bits(24, 8); + if (dest >= 16 || insn.words[1] || insn.words[2] != 0xc000 || native_predicate != 7) + return fail(index, "native warp sum mode"); + emit(0x50d8000000070000ULL | 250 | (u64{7} << 39) | (u64{7} << 45) | (u64{1} << 48)); + mov_imm(static_cast(dest) + 216, 0); + for (u32 lane = 0; lane < 32; ++lane) { + emit(0xef10000000070000ULL | 251 | (u64{source} << 8) | (u64{lane} << 20) | + (u64{1} << 28) | (u64{1} << 29) | (u64{31} << 34) | (u64{7} << 48)); + emit(0x3828000000070000ULL | 252 | (u64{250} << 8) | (u64{lane} << 20)); + emit(0x0400000000070000ULL | 252 | (u64{252} << 8) | (u64{1} << 20)); + emit(0x5c10000000070000ULL | 252 | (u64{255} << 8) | (u64{252} << 20) | (u64{1} << 48)); + emit(0x5c40000000070000ULL | 251 | (u64{251} << 8) | (u64{252} << 20) | (u64{7} << 48)); + emit(0x5c10000000070000ULL | (dest + 216) | (u64{dest + 216} << 8) | (u64{251} << 20)); + } + break; + } + case 0xf8c: { // ATOMS.POPC.INC.32 with no returned value: one increment per active lane. + const u32 stride = insn.Bits(78, 2); + if (dest != 255 || insn.words[1] || (insn.words[2] & ~0xc000U) != 0x0d00003f) + return fail(index, "native shared increment mode"); + u32 address = insn.Bits(24, 8); + if (stride && address != 255) { + emit(0x3848000000000000ULL | predicate | 251 | (u64{address} << 8) | + (u64{stride + 1} << 20)); + address = 251; + } + mov_imm(250, 1); + emit(0xec00000000000000ULL | predicate | 255 | (u64{address} << 8) | (u64{250} << 20)); + break; + } + case 0x38c: { // ATOMS.ADD.U32, naturally aligned shared byte offset. + const u32 offset = insn.Bits(40, 24); + if (insn.words[2] || (offset & 3)) return fail(index, "native shared atomic mode"); + emit(0xec00000000000000ULL | predicate | dest | + (u64{insn.Bits(24, 8)} << 8) | (u64{insn.Bits(32, 8)} << 20) | + (u64{offset >> 2} << 30)); + break; + } + case 0x9a8: { // ATOMG.E.ADD.STRONG.GPU U32 through a uniform pointer. + const u32 address = uniform_register(insn.Bits(64, 6)); + if (insn.Bits(24, 8) != 255 || insn.Bits(40, 24) || + (insn.words[2] & ~0x3fU) != 0x081ee140 || address == ~0U || (address & 1)) + return fail(index, "native global atomic mode"); + emit(0xed00000000000000ULL | predicate | dest | (u64{address} << 8) | + (u64{insn.Bits(32, 8)} << 20) | (u64{1} << 48)); + break; + } + case 0xb97: { // SULD.P.2D.IGN, direct descriptor, channel mask at 72-75. + const u32 binding = insn.Bits(54, 5), offset = insn.Bits(40, 14); + const u32 coordinate = insn.Bits(24, 8); + const u32 channels = insn.Bits(72, 4); + if ((insn.words[2] & ~0xf00U) != 0x001e0000 || channels == 0 || insn.Bits(61, 3) != 3 || + binding >= 18 || offset > 0x1fff || insn.Bits(59, 2) || dest == 255 || coordinate > 253) + return fail(index, "native formatted surface load mode"); + if (result.required_texture_buffer != ~0U && result.required_texture_buffer != binding) + return fail(index, "native surface descriptor buffer conflict"); + result.required_texture_buffer = binding; + emit(base(0xeb00000000000000ULL) | (u64{coordinate} << 8) | + (u64{channels} << 20) | (u64{3} << 33) | (u64{offset} << 36) | (u64{1} << 51)); + break; + } + case 0xb9b: { // SUST.P.2D/3D.IGN, direct descriptor and all four channels. + const u32 binding = insn.Bits(54, 5), offset = insn.Bits(40, 14); + const u32 data = insn.Bits(32, 8), coordinate = insn.Bits(24, 8); + const u32 dimension = insn.Bits(61, 3); + if (insn.words[2] != 0x00100f00 || (dimension != 3 && dimension != 5) || + binding >= 18 || offset > 0x1fff || insn.Bits(59, 2) || data > 252 || coordinate > (dimension == 5 ? 252U : 253U)) + return fail(index, "native formatted surface store mode"); + if (result.required_texture_buffer != ~0U && result.required_texture_buffer != binding) + return fail(index, "native surface descriptor buffer conflict"); + result.required_texture_buffer = binding; + emit(0xeb20000000000000ULL | predicate | data | (u64{coordinate} << 8) | + (u64{15} << 20) | (u64{dimension} << 33) | (u64{offset} << 36) | (u64{1} << 51)); + break; + } + case 0x301: // BREV register. + if (insn.words[2] || insn.words[1] > 255) return fail(index, "native bit reverse mode"); + emit(base(0x3828000000000000ULL) | (u64{insn.words[1]} << 8) | + (u64{1} << 39) | (u64{1} << 40)); + break; + case 0x219: { // U32 register shifts with a zero word outside the selected result. + const u32 low = insn.Bits(24, 8), high = insn.Bits(64, 8), shift = insn.Bits(32, 8); + const bool right = insn.Bits(76, 1) != 0, high_result = insn.Bits(80, 1) != 0; + if (insn.words[1] > 255 || insn.Bits(73, 2) != 3 || + (insn.words[2] & ~0x000118ffU) != 0x00000600U || + (right ? (!high_result || low != 255) : (high_result || high != 255))) + return fail(index, "native register shift mode"); + emit(base(right ? 0x5cf8000000000000ULL : 0x5bf8000000000000ULL) | + (u64{right ? high : 255U} << 8) | (u64{shift} << 20) | + (u64{right ? 255U : low} << 39) | (u64{insn.Bits(75, 1)} << 50)); + break; + } + case 0x300: case 0xd00: { // FLO, including a uniform source. + if ((insn.words[2] & ~(0x600U | (opcode == 0xd00 ? 0x08000000U : 0U))) != 0x000e0000U) + return fail(index, "native leading bit mode"); + const u32 source = operand(1, false); + emit(base(0x5c30000000000000ULL) | (u64{source} << 20) | + (u64{insn.Bits(63, 1)} << 40) | (u64{insn.Bits(74, 1)} << 41) | + (u64{insn.Bits(73, 1)} << 48)); + break; + } + case 0x309: case 0xd09: { // POPC, including a uniform source. + if (insn.words[2] != (opcode == 0xd09 ? 0x08000000U : 0U)) + return fail(index, "native population count mode"); + const u32 source = operand(1, false); + emit(base(0x5c08000000000000ULL) | (u64{source} << 20) | (u64{insn.Bits(63, 1)} << 40)); + break; + } + case 0x918: + if (insn.words[1] || insn.words[2]) return fail(index, "native NOP mode"); + emit(0x50b0000000000f00ULL | predicate); + break; + case 0x992: { // MEMBAR.SC: preserve workgroup/device visibility. + const u32 scope = insn.Bits(76, 3); + if (insn.words[1] || insn.words[2] != (scope << 12) || + (scope != 0 && scope != 2 && scope != 3 && scope != 5)) + return fail(index, "native memory barrier scope"); + const u32 legacy_scope = scope == 0 ? 0 : scope == 2 ? 1 : scope == 3 ? 2 : 3; + emit(0xef98000000000000ULL | predicate | (u64{legacy_scope} << 8)); + break; + } + case 0x9ab: // ERRBAR: wait for prior memory operations on the host backend. + if (insn.words[1] || insn.words[2]) return fail(index, "native error barrier mode"); + emit(0xef98000000000000ULL | predicate | (u64{1} << 8)); + break; + case 0xb1d: // BAR.SYNC 0, full workgroup. + if (insn.words[1] || insn.words[2] || native_predicate != 7) + return fail(index, "native workgroup barrier mode"); + emit(0xf0a8000000000000ULL | predicate | (u64{1} << 43) | + (u64{1} << 44) | (u64{1} << 39) | (u64{7} << 39)); + break; + case 0x388: case 0x984: { // Shared memory load/store with scaled register offsets. + const u32 width = insn.Bits(73, 3), stride = insn.Bits(78, 2); + if (width > 6 || insn.words[2] != ((width << 9) | (stride << 14)) || + (opcode == 0x984 && insn.Bits(32, 8))) + return fail(index, "native shared memory addressing mode"); + u32 address = insn.Bits(24, 8); + if (stride && address != 255) { + emit(0x3848000000000000ULL | predicate | 250 | + (u64{address} << 8) | (u64{stride + 1} << 20)); + address = 250; + } + const u32 data = opcode == 0x388 ? insn.Bits(32, 8) : static_cast(dest); + emit((opcode == 0x388 ? 0xef58000000000000ULL : 0xef48000000000000ULL) | + predicate | data | (u64{address} << 8) | (u64{insn.Bits(40, 24)} << 20) | + (u64{width} << 48)); + break; + } + case 0x919: case 0x9c3: { // S2R / S2UR thread and workgroup identities. + const u32 special = insn.Bits(72, 8); + if (insn.words[1] || insn.words[2] != (special << 8) || + (special != 0 && special != 33 && special != 34 && special != 35 && + special != 37 && special != 38 && special != 39 && + special != 56 && special != 57 && special != 58 && special != 59 && special != 60)) + return fail(index, "native special register"); + const u32 target = opcode == 0x9c3 ? uniform_register(static_cast(dest)) : static_cast(dest); + if (target == ~0U) return fail(index, "native uniform special destination"); + emit(0xf0c8000000000000ULL | predicate | target | (u64{special} << 20)); + break; + } + case 0x28c: case 0x88c: { // UISETP, signed (bit 73) or unsigned, with UPT accumulation. + const u32 output = insn.Bits(81, 3); + const u32 a = uniform_register(insn.Bits(24, 8)), b = opcode == 0x88c ? 252U : uniform_register(insn.Bits(32, 8)); + if (output > 1 || insn.Bits(84, 3) != 7 || insn.Bits(87, 4) != 7 || + (insn.Bits(72, 4) & ~2U) || insn.Bits(12, 4) != 7 || a == ~0U || b == ~0U) + return fail(index, "native uniform comparison mode"); + const u32 temporary = 6 - output, mask = 1U << temporary; + if (opcode == 0x88c) mov_imm(252, insn.words[1]); + emit(0x38e8000000070000ULL | 253 | (u64{255} << 8) | (u64{mask} << 20)); + emit(0x5b60000000070000ULL | 7 | (u64{temporary} << 3) | + (u64{a} << 8) | (u64{b} << 20) | (u64{7} << 39) | (u64{insn.Bits(73, 1)} << 48) | + (u64{insn.Bits(76, 3)} << 49)); + emit(0x38e8000000070000ULL | 239 | (u64{239} << 8) | (u64{mask} << 20)); + emit(0x38f0000000070000ULL | (u64{253} << 8) | (u64{mask} << 20)); + break; + } + case 0x287: case 0x887: { // USEL, UP0/UP1; preserve ordinary predicates. + const u32 condition = insn.Bits(87, 3); + const u32 target = uniform_register(static_cast(dest)); + const u32 a = uniform_register(insn.Bits(24, 8)), b = opcode == 0x887 ? 252U : uniform_register(insn.Bits(32, 8)); + if (condition > 1 || target == ~0U || a == ~0U || b == ~0U || + insn.Bits(12, 4) != 7 || (insn.words[2] & ~0x07800000U) != 0x08000000U) + return fail(index, "native uniform select mode"); + if (opcode == 0x887 && index > 0 && condition == 0 && insn.words[1] <= 0x7ffff && + code[index - 1].Bits(0, 12) == 0x88c && + code[index - 1].Bits(24, 8) == insn.Bits(24, 8) && + code[index - 1].words[1] == insn.words[1] && code[index - 1].words[2] == 0x0bf01070 && + insn.words[2] == 0x08000000) { + // Preserve the unsigned clamp idiom for indirect-table analysis. + emit(0x3820000000070000ULL | target | (u64{a} << 8) | + (u64{insn.words[1]} << 20) | (u64{7} << 39)); + break; + } + if (opcode == 0x887) mov_imm(252, insn.words[1]); + const u32 temporary = 6 - condition, mask = 1U << temporary; + emit(0x38e8000000070000ULL | 253 | (u64{255} << 8) | (u64{mask} << 20)); + emit(0x38f0000000070000ULL | (u64{239} << 8) | (u64{mask} << 20)); + emit(0x5ca0000000070000ULL | target | (u64{a} << 8) | (u64{b} << 20) | + (u64{temporary} << 39) | (u64{insn.Bits(90, 1)} << 42)); + emit(0x38f0000000070000ULL | (u64{253} << 8) | (u64{mask} << 20)); + break; + } + case 0x386: case 0x986: { // Global store of 32/64/128 bits, 64-bit address. + const u32 data = insn.Bits(32, 8); + const u32 size = insn.Bits(73, 3); + if (size < 4 || size > 6 || (size == 5 && (data & 1)) || (size == 6 && (data & 3)) || + data + (1U << (size - 4)) > 246) + return fail(index, "native global store size"); + u32 address = insn.Bits(24, 8); + bool extended = insn.Bits(72, 1) != 0; + if (opcode == 0x986) { + const u32 uniform_address = uniform_register(insn.Bits(64, 6)); + if ((insn.words[2] & ~(0x04000000U | 0x3fU | 0xe00U)) != 0x0810a100 || + uniform_address == ~0U || (uniform_address & 1) || !extended) + return fail(index, "native global store uniform addressing mode"); + if (address == 255) address = uniform_address; + else { + if (insn.Bits(90, 1) && ((address & 1) || address >= 254)) + return fail(index, "native global store register address pair"); + emit(0x5c10000000000000ULL | predicate | 250 | (u64{uniform_address} << 8) | + (u64{address} << 20) | (u64{1} << 47)); + emit(0x5c10000000000000ULL | predicate | 251 | (u64{uniform_address + 1} << 8) | + (u64{insn.Bits(90, 1) ? address + 1 : 255} << 20) | (u64{1} << 43)); + address = 250; + } + } else if ((insn.words[2] & ~0xe00U) != 0x0010a100) { + return fail(index, "native global store register addressing mode"); + } + emit(0xeed8000000000000ULL | predicate | data | + (u64{address} << 8) | (u64{insn.Bits(40, 24)} << 20) | + (u64{extended} << 45) | (u64{size} << 48)); + break; + } + case 0x95b: // Fragment discard; execution predicate is preserved. + if (insn.words[1] || insn.words[2] != 0x03800000) + return fail(index, "native fragment discard mode"); + emit(0xe33000000000000fULL | predicate); + break; + case 0x305: case 0xb05: { // F2I NTZ, register/constant-buffer source + const bool half = opcode == 0x305 && insn.words[2] == 0x0010f000; + const bool small = insn.words[2] == 0x00212800; + const bool floor = (insn.words[2] & ~0x100U) == 0x00207000; + const bool nearest = (insn.words[2] & ~0x100U) == 0x00203000; + if ((!half && !small && !floor && !nearest && (insn.words[2] & ~0x100U) != 0x0021f000) || (opcode == 0x305 && (insn.words[1] & 0xffffff00U))) { + return fail(index, "native float-to-integer conversion mode"); + } + const bool is_signed = insn.Bits(72, 1); + u32 source = opcode == 0xb05 ? materialize(2, 0, 252, false, false) : insn.Bits(32, 8); + if (half) { + extract_half({source, 2, false, false}, 0, 252); + source = 252; + } + if (!is_signed) { + // Clear NaNs without consuming a guest predicate. Clamp the + // conversion below 2^32, retaining an overflow mask for saturation. + emit(0x5800000000000000ULL | predicate | 250 | (u64{source} << 8) | + (u64{source} << 20) | (u64{7} << 39) | (u64{7} << 48)); + emit(0x5c40000000000000ULL | predicate | 250 | (u64{source} << 8) | + (u64{250} << 20) | (u64{7} << 48)); + if (!small) { + mov_imm(252, 0x4f800000); + emit(0x5800000000000000ULL | predicate | 251 | (u64{250} << 8) | + (u64{252} << 20) | (u64{7} << 39) | (u64{6} << 48)); + mov_imm(252, 0x4f7fffff); + emit(0x5c60000000000000ULL | predicate | 250 | (u64{250} << 8) | + (u64{252} << 20) | (u64{7} << 39)); + } + source = 250; + } + emit(base(0x5cb0000000000000ULL) | (u64{small ? 1U : 2U} << 8) | (u64{2} << 10) | + (u64{is_signed} << 12) | (u64{source} << 20) | + (u64{small || nearest ? 0U : floor ? 1U : 3U} << 39) | (u64{insn.Bits(80, 1)} << 44)); + if (!is_signed && !small) { + emit(base(0x5c40000000000000ULL) | (dest << 8) | (u64{251} << 20) | + (u64{1} << 41) | (u64{7} << 48)); + } + break; + } + case 0xb6f: { // TXQ texture dimensions, two packed scalar outputs. + if (insn.Bits(59, 5) || insn.Bits(64, 8) != 255 || insn.Bits(72, 4) != 3 || + insn.Bits(76, 14) || insn.Bits(91, 1) || insn.Bits(54, 5) != 2 || + insn.Bits(40, 14) > 0x1fff || (dest & 1) || dest >= 254) + return fail(index, "native texture dimension query mode"); + result.required_texture_buffer = 2; + emit(0xdf48000000000000ULL | predicate | dest | (u64{insn.Bits(24, 8)} << 8) | + (u64{1} << 22) | (u64{3} << 31) | (u64{insn.Bits(40, 14)} << 36) | + (u64{insn.Bits(90, 1)} << 49)); + break; + } + case 0xb66: { // TLD.SCR.LZ, direct texel fetch + const u32 mask = insn.Bits(72, 4); + const u32 dimension = insn.Bits(61, 3); + const bool array = dimension == 5; + const bool volume = dimension == 2; + u32 source_a = insn.Bits(24, 8), source_b = insn.Bits(32, 8); + const u32 second_dest = insn.Bits(64, 8); + const bool split = second_dest != 255; + constexpr std::array rg_masks{1, 2, 4, 8, 3, 9, 10, 12}; + constexpr std::array rgba_masks{7, 11, 13, 14, 15}; + u32 swizzle = ~0U; + if (split) { + for (u32 i = 0; i < rgba_masks.size(); ++i) if (rgba_masks[i] == mask) swizzle = i; + } else { + for (u32 i = 0; i < rg_masks.size(); ++i) if (rg_masks[i] == mask) swizzle = i; + } + if (!insn.Bits(60, 1) || (dimension != 0 && dimension != 1 && !array && !volume) || + swizzle == ~0U || insn.Bits(76, 5) || + insn.Bits(81, 3) != 7 || insn.Bits(87, 3) != 1 || + insn.Bits(59, 1) || + insn.Bits(40, 14) > 0x1fff || insn.Bits(54, 5) >= 18 || + dest == 255 || (mask & (mask - 1) && ((dest & 1) || dest >= 254)) || + (split && mask == 15 && ((second_dest & 1) || second_dest >= 254)) || + (array && ((source_a & 1) || source_a >= 249)) || + (volume && ((source_a & 1) || source_a >= 254))) { + return fail(index, "native texel fetch mode or descriptor binding"); + } + const u32 binding = insn.Bits(54, 5); + if (result.required_texture_buffer != ~0U && result.required_texture_buffer != binding) + return fail(index, "native texture descriptor buffer conflict"); + result.required_texture_buffer = binding; + if (array) { + emit(0x5c98000000000000ULL | predicate | 250 | (u64{15} << 39) | (u64{source_a + 1} << 20)); + emit(0x5c98000000000000ULL | predicate | 251 | (u64{15} << 39) | (u64{source_b} << 20)); + source_b = 250; + } + emit(0xda00000000000000ULL | predicate | dest | (u64{source_a} << 8) | + (u64{source_b} << 20) | (u64{insn.Bits(64, 8)} << 28) | + (u64{insn.Bits(40, 14)} << 36) | (u64{insn.Bits(90, 1)} << 49) | + (u64{swizzle} << 50) | (u64{array ? 8U : volume ? 7U : dimension == 0 ? 0U : 2U} << 53)); + break; + } + case 0x91a: // DEPBAR.LE SB5, 1: scoreboard wait, represented by IR data dependencies. + if (insn.words[0] != 0x791a || (insn.words[1] & ~0xfc0U) != 0xd000 || insn.words[2]) { + return fail(index, "native dependency barrier mode"); + } + break; + case 0xb60: { // TEX, scalar 2D implicit-LOD sample + const u32 mask = insn.Bits(72, 4); + const u32 second_dest = insn.Bits(64, 8); + constexpr std::array rg_masks{1, 2, 4, 8, 3, 9, 10, 12}; + constexpr std::array rgba_masks{7, 11, 13, 14, 15}; + u32 swizzle = ~0U; + const bool split = second_dest != 255; + const bool depth_compare = insn.Bits(78, 1); + const u32 dimension = insn.Bits(61, 3), lod = insn.Bits(87, 3); + const u32 source_a = insn.Bits(24, 8), source_b = insn.Bits(32, 8); + if (dimension == 5 && depth_compare && lod == 0 && !split && mask == 1) { + if (!insn.Bits(60, 1) || insn.Bits(76, 2) || insn.Bits(81, 3) != 7 || + insn.Bits(59, 1) || insn.Bits(40, 14) > 0x1fff || insn.Bits(54, 5) != 2 || + (source_a & 1) || (source_b & 1) || source_a >= 248 || source_b >= 248 || dest >= 248) + return fail(index, "native array depth comparison sample mode"); + result.required_texture_buffer = 2; + for (u32 i = 0; i < 3; ++i) { + const u32 source = i < 2 ? source_a + i : source_b; + emit(0x5c98000000000000ULL | predicate | (250 + i) | + (u64{15} << 39) | (u64{source} << 20)); + } + emit(0xc000000000000000ULL | predicate | dest | (u64{250} << 8) | + (u64{source_b + 1} << 20) | (u64{3} << 28) | (u64{1} << 31) | + (u64{insn.Bits(40, 14)} << 36) | (u64{insn.Bits(90, 1)} << 49) | + (u64{1} << 50) | (u64{7} << 51)); + break; + } + // General form: gather the native operand list (array index first, coordinates, + // then LOD/bias and the depth reference) into Maxwell TEX's coordinate and meta + // registers, sample into scratch and copy the packed components out. Native packs + // up to four values as halves of Ra and Rb (1+1, 2+1, 2+2), five or more as 4+n. + const auto general_tex = [&]() -> bool { + constexpr std::array coordinate_count{1, 2, 3, 3, 2, 3, 0, 4}; + constexpr std::array maxwell_type{0, 2, 4, 6, 1, 3, 0, 7}; + if (dimension == 6 || insn.Bits(76, 2) || insn.Bits(81, 3) != 7 || insn.Bits(59, 1) || + insn.Bits(40, 14) > 0x1fff || insn.Bits(54, 5) >= 18 || (lod > 3 && lod != 6 && lod != 7) || + (depth_compare && mask != 1) || mask == 0) + return false; + const u32 coordinates = coordinate_count[dimension]; + const bool has_lod = lod == 2 || lod == 3 || lod == 6 || lod == 7; + const u32 count = coordinates + (has_lod ? 1 : 0) + (depth_compare ? 1 : 0); + const u32 in_a = count <= 4 ? (count + 1) / 2 : 4; + const auto value = [&](u32 i) -> u32 { + if (i < in_a) return source_a == 255 ? 255U : source_a + i; + return source_b == 255 ? 255U : source_b + (i - in_a); + }; + if ((source_a != 255 && source_a + in_a > 246) || + (count > in_a && source_b != 255 && source_b + (count - in_a) > 246)) + return false; + const u32 binding = insn.Bits(54, 5); + if (result.required_texture_buffer != ~0U && result.required_texture_buffer != binding) + return false; + result.required_texture_buffer = binding; + const auto copy = [&](u32 target, u32 source) { + emit(0x5c98000000000000ULL | predicate | target | (u64{15} << 39) | (u64{source} << 20)); + }; + // Coordinates to R250.., LOD/bias then reference to R246/R247. + for (u32 i = 0; i < coordinates; ++i) copy(250 + i, value(i)); + u32 meta = 246; + if (has_lod) copy(meta++, value(coordinates)); + if (depth_compare) copy(meta++, value(coordinates + (has_lod ? 1 : 0))); + emit(0xc000000000000000ULL | predicate | 250 | (u64{250} << 8) | + (u64{meta == 246 ? 255U : 246U} << 20) | (u64{maxwell_type[dimension]} << 28) | + (u64{mask} << 31) | (u64{insn.Bits(40, 14)} << 36) | + (u64{insn.Bits(90, 1)} << 49) | (u64{depth_compare} << 50) | (u64{7} << 51) | + (u64{lod} << 55)); + const u32 components = static_cast(std::popcount(mask)); + for (u32 i = 0; i < components; ++i) { + u32 target = 255; + if (!split) target = dest == 255 ? 255U : static_cast(dest) + i; + else if (i < 2) target = dest == 255 ? 255U : static_cast(dest) + i; + else target = second_dest + (i - 2); + if (target != 255) copy(target, 250 + i); + } + return true; + }; + if (split) { + for (u32 i = 0; i < rgba_masks.size(); ++i) if (rgba_masks[i] == mask) swizzle = i; + } else { + for (u32 i = 0; i < rg_masks.size(); ++i) if (rg_masks[i] == mask) swizzle = i; + } + // SM70 scalar coordinates and split RG/BA destinations match TEXS. + // Descriptor offsets are words in both encodings, but only native + // explicitly encodes the constant-buffer binding. + if (!insn.Bits(60, 1) || (dimension != 1 && dimension != 5) || + swizzle == ~0U || insn.Bits(76, 2) || + insn.Bits(81, 3) != 7 || (lod > 1 && (lod != 3 || dimension != 1 || depth_compare)) || + (lod == 3 && ((source_a & 1) || source_a >= 249)) || + insn.Bits(59, 1) || + insn.Bits(40, 14) > 0x1fff || insn.Bits(54, 5) != 2 || + dest == 255 || (mask & (mask - 1) && ((dest & 1) || dest >= 254)) || + (split && mask == 15 && ((second_dest & 1) || second_dest >= 254)) || + (depth_compare && (dimension != 1 || split || mask != 1 || (source_a & 1) || source_a >= 249)) || + (dimension == 5 && ((source_a & 1) || source_a >= 249))) { + if (!general_tex()) return fail(index, "native texture mode or descriptor binding"); + break; + } + result.required_texture_buffer = 2; + emit(0xd800000000000000ULL | predicate | dest | (u64{insn.Bits(24, 8)} << 8) | + (u64{insn.Bits(32, 8)} << 20) | (u64{insn.Bits(64, 8)} << 28) | + (u64{insn.Bits(40, 14)} << 36) | (u64{insn.Bits(90, 1)} << 49) | + (u64{swizzle} << 50) | (u64{dimension == 5 ? (lod ? 8U : 7U) : depth_compare ? (lod ? 6U : 4U) : (lod == 3 ? 3U : lod ? 2U : 1U)} << 53)); + break; + } + case 0x816: { // Captured PRMT half-word concatenation + if ((insn.words[1] != 0x5410 && insn.words[1] != 0x5432 && insn.words[1] != 0x7610 && insn.words[1] != 0x7632) || + (insn.words[2] & ~0xffU)) return fail(index, "native permutation mode"); + const u32 a = insn.Bits(24, 8), b = insn.Bits(64, 8); + if (insn.words[1] != 0x5432 && insn.words[1] != 0x7632) { + emit(0x0400000000000000ULL | predicate | 250 | (u64{a} << 8) | (u64{0xffff} << 20)); + } else { + emit(0x3828000000000000ULL | predicate | 250 | (u64{a} << 8) | (u64{16} << 20)); + } + if (insn.words[1] == 0x7610 || insn.words[1] == 0x7632) { + emit(0x0400000000000000ULL | predicate | 251 | (u64{b} << 8) | (u64{0xffff0000} << 20)); + emit(base(0x5c10000000000000ULL) | (u64{250} << 8) | (u64{251} << 20)); + } else { + emit(base(0x5c18000000000000ULL) | (u64{b} << 8) | (u64{250} << 20) | (u64{16} << 39)); + } + break; + } + case 0x307: { // FRND.F32, round to an integral floating-point value + if (insn.Bits(75, 2) != 2 || insn.Bits(84, 2) != 2 || insn.Bits(81, 3) || + insn.Bits(86, 6)) return fail(index, "native floating round format"); + emit(base(0x5ca8000000000000ULL) | (u64{2} << 8) | (u64{2} << 10) | + (u64{insn.Bits(32, 8)} << 20) | (u64{8 + insn.Bits(78, 2)} << 39) | + (u64{insn.Bits(80, 1)} << 44) | (u64{insn.Bits(63, 1)} << 45) | + (u64{insn.Bits(62, 1)} << 49)); + break; + } + case 0x321: // ALD + case 0x322: { // AST + if (insn.Bits(76, 2) || insn.Bits(79, 1)) return fail(index, "native indexed/patch attribute access"); + const u32 address = insn.Bits(40, 10); + const u32 components = insn.Bits(74, 2) + 1; + if (address % 4 || address + components * 4 > 0x400 || insn.Bits(24, 8) != 255) { + return fail(index, "native attribute address or indexing"); + } + const u32 value_reg = opcode == 0x321 ? static_cast(dest) : insn.Bits(32, 8); + if (value_reg + components > 255) return fail(index, "native attribute register range"); + const u32 vertex = opcode == 0x321 ? insn.Bits(32, 8) : insn.Bits(64, 8); + const u64 translated = (opcode == 0x321 ? 0xefd8000000000000ULL : 0xeff0000000000000ULL) | + predicate | value_reg | (u64{255} << 8) | (u64{address} << 20) | + (u64{vertex} << 39) | (u64{components - 1} << 47); + result.instructions.push_back(translated); + break; + } + case 0x387: case 0x983: { // Captured scalar local spill at byte offset zero. + const bool store = opcode == 0x387; + if (insn.Bits(24, 8) != 255 || (store && (insn.Bits(16, 8) || insn.words[1] > 247)) || + (!store && (insn.words[1] || dest >= 248)) || + insn.words[2] != (store ? 0x00100800U : 0x00300800U)) + return fail(index, "native local scalar spill mode"); + emit((store ? 0xef50000000000000ULL : 0xef40000000000000ULL) | predicate | + (store ? insn.words[1] : dest) | (u64{255} << 8) | (u64{4} << 48)); + break; + } + case 0x803: { // P2R PR, register merge source, immediate mask, byte select (76-77) + if (insn.words[1] > 0x7f || (insn.words[2] & ~0x3000U)) + return fail(index, "native predicate-to-register mode"); + emit(base(0x38e8000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{insn.words[1]} << 20) | (u64{insn.Bits(76, 2)} << 41)); + break; + } + case 0x804: { // R2P PR, register, immediate mask, byte select (76-77) + if (insn.Bits(16, 8) || insn.words[1] > 0x7f || (insn.words[2] & ~0x3000U)) + return fail(index, "native register-to-predicate mode"); + emit(0x38f0000000000000ULL | predicate | (u64{insn.Bits(24, 8)} << 8) | + (u64{insn.words[1]} << 20) | (u64{insn.Bits(76, 2)} << 41)); + break; + } + case 0xc02: // MOV uniform register + case 0x202: // MOV register + case 0x802: // MOV immediate + case 0xa02: { // MOV constant buffer + if (insn.Bits(72, 4) != 15 || (insn.Bits(76, 16) & ~(opcode == 0xc02 ? 0x8000U : 0U))) return fail(index, "native MOV modifiers"); + if (opcode == 0x802) { + result.instructions.push_back(base(0x0100000000000000ULL) | (u64{15} << 12) | + (u64{insn.words[1]} << 20)); + } else if (opcode == 0xc02) { + if (insn.words[1] >= 16 || insn.words[2] != 0x08000f00) return fail(index, "native uniform move source"); + emit(base(0x5c98000000000000ULL) | (u64{15} << 39) | (u64{insn.words[1] + 216} << 20)); + } else if (opcode == 0x202) { + if (insn.Bits(40, 24)) return fail(index, "native MOV register modifiers"); + result.instructions.push_back(base(0x5c98000000000000ULL) | (u64{15} << 39) | + (u64{insn.Bits(32, 8)} << 20)); + } else { + if (!valid_cbuf() || insn.Bits(62, 2)) return fail(index, "native MOV constant buffer"); + result.instructions.push_back(base(0x4c98000000000000ULL) | (u64{15} << 39) | cbuf()); + } + break; + } + case 0x223: case 0x423: case 0x623: case 0x823: case 0xa23: case 0xc23: case 0xe23: { // FFMA + if (insn.Bits(76, 1) || (insn.Bits(81, 11) & ~(opcode == 0xe23 || opcode == 0xc23 ? 0x400U : 0U))) return fail(index, "native FFMA denormal/source mode"); + const u32 a = operand(0), b = operand(1), c = operand(2); + emit(base(0x5980000000000000ULL) | (u64{a} << 8) | (u64{b} << 20) | (u64{c} << 39) | + (u64{insn.Bits(77, 1)} << 50) | (u64{insn.Bits(78, 2)} << 51) | (u64{insn.Bits(80, 1)} << 53)); + break; + } + case 0x220: case 0x820: case 0xa20: case 0xc20: { // FMUL + const u32 scale = insn.Bits(84, 3); + if (insn.Bits(76, 1) || scale == 0 || insn.Bits(81, 3) || + insn.Bits(87, 4) || insn.Bits(91, 1) != (opcode == 0xc20) || + (scale != 4 && (!insn.Bits(80, 1) || insn.Bits(78, 2)))) + return fail(index, "native FMUL scale/mode"); + const u32 a = operand(0), b = operand(1); + emit(base(0x5c68000000000000ULL) | (u64{a} << 8) | (u64{b} << 20) | + (u64{insn.Bits(78, 2)} << 39) | (u64{std::array{0, 3, 2, 1, 0, 6, 5, 4}[scale]} << 41) | (u64{insn.Bits(80, 1)} << 44) | (u64{insn.Bits(77, 1)} << 50)); + break; + } + case 0x221: case 0x421: case 0x621: case 0x821: case 0xa21: case 0xe21: { // FADD + if (insn.Bits(81, 10) || insn.Bits(91, 1) != (opcode == 0xe21)) return fail(index, "native FADD mode"); + const u32 a = operand(0), b = operand(insn.Bits(9, 3) == 2 || insn.Bits(9, 3) == 3 || insn.Bits(9, 3) == 7 ? 2 : 1); + emit(base(0x5c58000000000000ULL) | (u64{a} << 8) | (u64{b} << 20) | + (u64{insn.Bits(78, 2)} << 39) | (u64{insn.Bits(80, 1)} << 44) | (u64{insn.Bits(77, 1)} << 50)); + break; + } + case 0x209: case 0x809: case 0xa09: { // FMNMX + if (insn.Bits(74, 6) || insn.Bits(81, 6) || insn.Bits(91, 1)) + return fail(index, "native floating min/max mode"); + const u32 a = operand(0), b = operand(1); + emit(base(0x5c60000000000000ULL) | (u64{a} << 8) | (u64{b} << 20) | + (u64{insn.Bits(87, 3)} << 39) | (u64{insn.Bits(90, 1)} << 42) | + (u64{insn.Bits(80, 1)} << 44)); + break; + } + case 0x308: case 0xb08: { // MUFU + const u32 operation = insn.Bits(74, 6); + if (operation > 8 || operation == 6 || operation == 7 || insn.Bits(80, 12)) return fail(index, "native MUFU mode"); + if (insn.Bits(73, 1)) { + if (opcode != 0x308 || insn.Bits(60, 2) > 1) + return fail(index, "native half special-function source"); + emit(0x5ca8000000000000ULL | predicate | 250 | (u64{2} << 8) | (u64{1} << 10) | + (u64{insn.Bits(32, 8)} << 20) | (u64{insn.Bits(60, 1)} << 41) | + (u64{insn.Bits(63, 1)} << 45) | (u64{insn.Bits(62, 1)} << 49)); + emit(0x5080000000000000ULL | predicate | 250 | (u64{250} << 8) | (u64{operation} << 20)); + emit(base(0x5ca8000000000000ULL) | (u64{1} << 8) | (u64{2} << 10) | (u64{250} << 20)); + } else { + const u32 src = operand(1); + emit(base(0x5080000000000000ULL) | (u64{src} << 8) | (u64{operation} << 20)); + } + break; + } + case 0x526: { // IPA.PASS / CONSTANT with the attribute indexed by a uniform register + const u32 index_register = uniform_register(insn.Bits(24, 8)); + if (index_register == ~0U || insn.Bits(32, 8) != 255 || + (insn.Bits(76, 4) != 0 && insn.Bits(76, 4) != 4) || insn.Bits(81, 3) != 7) + return fail(index, "native indexed IPA mode"); + if (index_register == 255) { + emit(base(0xe000000000000000ULL) | (u64{255} << 8) | (u64{255} << 20) | + (u64{insn.Bits(64, 8)} << 30) | (u64{insn.Bits(78, 1)} << 55)); + break; + } + // Maxwell IPA.IDX reads the whole byte address from the index register. + emit(0x1c00000000000000ULL | predicate | 250 | (u64{index_register} << 8) | + (u64{insn.Bits(64, 8) * 4} << 20)); + emit(base(0xe000000000000000ULL) | (u64{250} << 8) | (u64{255} << 20) | + (u64{insn.Bits(64, 8)} << 30) | (u64{1} << 38) | (u64{insn.Bits(78, 1)} << 55)); + break; + } + case 0x326: { // IPA.PASS / CONSTANT, default sample location + if (insn.Bits(32, 8) != 255 || (insn.Bits(76, 4) != 0 && insn.Bits(76, 4) != 4) || insn.Bits(81, 3) != 7) return fail(index, "native IPA interpolation mode"); + emit(base(0xe000000000000000ULL) | (u64{255} << 8) | (u64{255} << 20) | (u64{insn.Bits(64, 8)} << 30) | (u64{insn.Bits(78, 1)} << 55)); + break; + } + case 0x805: { // CS2R SRZ + if (insn.Bits(72, 8) != 255 || !insn.Bits(80, 1) || dest >= 249) return fail(index, "native special-register read"); + mov_imm(static_cast(dest), 0); + mov_imm(static_cast(dest) + 1, 0); + break; + } + case 0x306: { // I2F.F32.U8, low byte, round-to-nearest + if (insn.words[2] != 0x1000 || (insn.words[1] & 0xffffff00U)) + return fail(index, "native byte-to-float conversion mode"); + emit(base(0x5cb8000000000000ULL) | (u64{2} << 8) | (u64{insn.Bits(32, 8)} << 20)); + break; + } + case 0x245: case 0xa45: case 0xc45: { // I2FP.F32.U32 + if ((insn.words[2] & ~(0x400U | (opcode == 0xc45 ? 0x08000000U : 0U))) != 0x00201000) return fail(index, "native packed integer-to-float format"); + const u32 src = operand(1, false); + emit(base(0x5cb8000000000000ULL) | (u64{2} << 8) | (u64{2} << 10) | (u64{insn.Bits(74, 1)} << 13) | (u64{src} << 20)); + break; + } + case 0x207: case 0x807: case 0xa07: + case 0x208: case 0x808: case 0xa08: { // Bitwise SEL / FSEL + const u32 a = operand(0, false), b = operand(1, false); + emit(base(0x5ca0000000000000ULL) | (u64{a} << 8) | (u64{b} << 20) | + (u64{insn.Bits(87, 3)} << 39) | (u64{insn.Bits(90, 1)} << 42)); + break; + } + case 0x20a: case 0x80a: case 0xa0a: { // FSET.BF with PT/AND + if (insn.Bits(64, 8) || insn.Bits(74, 2) || insn.Bits(81, 6) || + insn.Bits(87, 4) != 7 || insn.Bits(91, 1)) + return fail(index, "native floating comparison output mode"); + const u32 a = operand(0), b = operand(1); + emit(base(0x5800000000000000ULL) | (u64{a} << 8) | (u64{b} << 20) | + (u64{7} << 39) | (u64{insn.Bits(76, 4)} << 48) | (u64{1} << 52) | + (u64{insn.Bits(80, 1)} << 55)); + break; + } + case 0x20b: case 0x80b: case 0xa0b: // FSETP + case 0x20c: case 0x80c: case 0xa0c: case 0xc0c: { // ISETP + const bool floating = (opcode & 0x1ff) == 0xb; + if (!floating && insn.Bits(72, 1)) return fail(index, "native extended integer comparison"); + const u32 a = operand(0, floating), b = operand(1, floating); + const u32 combine = insn.Bits(74, 2); + if (combine > 2) return fail(index, "native predicate combination"); + u64 out = (floating ? 0x5bb0000000000000ULL : 0x5b60000000000000ULL) | predicate | + insn.Bits(84, 3) | (u64{insn.Bits(81, 3)} << 3) | (u64{a} << 8) | (u64{b} << 20) | + (u64{insn.Bits(87, 3)} << 39) | (u64{insn.Bits(90, 1)} << 42) | (u64{combine} << 45); + if (floating) out |= (u64{insn.Bits(80, 1)} << 47) | (u64{insn.Bits(76, 4)} << 48); + else out |= (u64{insn.Bits(73, 1)} << 48) | (u64{insn.Bits(76, 3)} << 49); + emit(out); + break; + } + case 0x212: case 0x812: case 0xc12: { // Bounded LOP3 bitwise forms, no predicate output + const u32 lut = insn.Bits(72, 8); + if ((lut != 0xc0 && lut != 0xfc && lut != 0xf8 && lut != 0x48 && lut != 0xe2) || insn.Bits(80, 1) || + insn.Bits(87, 4) != 15 || + (opcode == 0x212 && insn.Bits(40, 24))) return fail(index, "native LUT/predicate mode"); + u32 b = opcode == 0xc12 ? uniform_register(insn.Bits(32, 8)) : insn.Bits(32, 8); + if (b == ~0U) return fail(index, "native uniform logic source"); + if (opcode == 0x812) { + mov_imm(251, insn.words[1]); + b = 251; + } + const u32 a = insn.Bits(24, 8), c = insn.Bits(64, 8); + const auto logical = [&](u32 target, u32 lhs, u32 rhs, u32 operation, u32 pred_target = 7) { + emit(0x5c40000000000000ULL | predicate | target | (u64{lhs} << 8) | + (u64{rhs} << 20) | (u64{operation} << 41) | + (u64{pred_target == 7 ? 0U : 3U} << 44) | (u64{pred_target} << 48)); + }; + const u32 pred_target = insn.Bits(81, 3); + if (lut == 0xe2) { + logical(250, a, b, 0); + logical(252, c, b, 0); + result.instructions.back() |= u64{1} << 40; + logical(static_cast(dest), 250, 252, 1, pred_target); + } else if (lut == 0xf8) { + logical(250, b, c, false); + logical(static_cast(dest), a, 250, 1, pred_target); + } else if (lut == 0x48) { + logical(250, a, c, 2); + logical(static_cast(dest), 250, b, 0, pred_target); + } else { + logical(static_cast(dest), a, b, lut == 0xfc ? 1U : 0U, pred_target); + } + break; + } + case 0x217: case 0x817: case 0xa17: case 0xc17: { // IMNMX register, immediate, cbuf or uniform operand. + if ((insn.words[2] & ~(0x04000200U | (opcode == 0xc17 ? 0x08000000U : 0U))) != 0x03800000 || + (opcode != 0xa17 && insn.words[1] > (opcode == 0x817 ? 0x7ffffU : 255U)) || + (opcode == 0xa17 && (!valid_cbuf() || insn.Bits(62, 2)))) + return fail(index, "native integer min/max mode"); + const u32 source = opcode == 0xc17 ? uniform_register(insn.words[1]) + : opcode == 0xa17 ? operand(1, false) + : insn.words[1]; + if (source == ~0U) return fail(index, "native uniform min/max operand"); + emit(base(opcode == 0x817 ? 0x3820000000000000ULL : 0x5c20000000000000ULL) | + (u64{insn.Bits(24, 8)} << 8) | (u64{source} << 20) | + (u64{insn.Bits(87, 4)} << 39) | (u64{insn.Bits(73, 1)} << 48)); + break; + } + case 0x819: { // SHF.HI, immediate shift + const bool is_signed = insn.Bits(73, 2) == 2; + if (!insn.Bits(80, 1) && !insn.Bits(76, 1) && insn.Bits(64, 8) == 255 && + insn.Bits(73, 2) == 3 && !insn.Bits(75, 1) && insn.words[1] <= 32) { + emit(base(0x3848000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | (u64{insn.words[1]} << 20)); + break; + } + if ((insn.Bits(73, 2) != 3 && !is_signed) || insn.Bits(75, 1) || !insn.Bits(80, 1) || insn.words[1] > 32) return fail(index, "native funnel-shift mode"); + const u32 low = insn.Bits(24, 8), high = insn.Bits(64, 8), shift = insn.words[1]; + if (insn.Bits(76, 1)) { + if (low != 255) return fail(index, "native nonzero low right-shift word"); + emit(base(0x3828000000000000ULL) | (u64{high} << 8) | (u64{shift} << 20) | + (u64{is_signed} << 48)); + } else { + if (is_signed) return fail(index, "native signed left funnel shift"); + emit(base(0x36f8000000000000ULL) | (u64{low} << 8) | (u64{high} << 39) | (u64{shift} << 20)); + } + break; + } + case 0xa11: { // LEA with constant base and carry predicate. + if (insn.Bits(72, 1) || insn.Bits(74, 1) || insn.Bits(80, 1) || + insn.Bits(81, 3) == 7 || insn.Bits(62, 2)) + return fail(index, "native constant LEA carry mode"); + const u32 b = operand(1, false); + emit(base(0x5c18000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{b} << 20) | (u64{insn.Bits(75, 5)} << 39)); + emit(0x5b60000000000000ULL | predicate | 7 | (u64{insn.Bits(81, 3)} << 3) | + (dest << 8) | (u64{b} << 20) | (u64{7} << 39) | (u64{1} << 49)); + break; + } + case 0x211: case 0x811: { // LEA low word, no carry result + if (opcode == 0x211 && insn.Bits(80, 1)) { + if (insn.Bits(72, 1) || insn.Bits(74, 1) || insn.Bits(81, 3) != 7 || + insn.Bits(32, 8) != 255 || insn.Bits(64, 8) != 255 || insn.Bits(62, 2)) + return fail(index, "native high effective address mode"); + emit(base(0x3828000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{32 - insn.Bits(75, 5)} << 20)); + break; + } + if (insn.Bits(72, 1) || insn.Bits(74, 1) || insn.Bits(80, 1) || insn.Bits(81, 3) != 7 || + (opcode == 0x211 && insn.Bits(62, 2))) return fail(index, "native LEA carry/high mode"); + u32 b = insn.Bits(32, 8); + if (opcode == 0x811) { mov_imm(251, insn.words[1]); b = 251; } + emit(base(0x5c18000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{b} << 20) | (u64{insn.Bits(75, 5)} << 39)); + break; + } + case 0x210: case 0x810: case 0xa10: case 0xc10: { // IADD3 with zero third source and no carry outputs + if (opcode == 0xa10 && insn.Bits(81, 3) != 7) { + if (insn.Bits(64, 8) != 255 || insn.Bits(73, 4) || insn.Bits(84, 7) != 0x7f || + insn.Bits(77, 4) != 15 || insn.Bits(63, 1) || insn.Bits(72, 1)) + return fail(index, "native constant add carry mode"); + const u32 b = operand(1, false); + emit(base(0x5c10000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | (u64{b} << 20)); + emit(0x5b60000000000000ULL | predicate | 7 | (u64{insn.Bits(81, 3)} << 3) | + (dest << 8) | (u64{b} << 20) | (u64{7} << 39) | (u64{1} << 49)); + break; + } + if (opcode == 0xa10 && insn.Bits(74, 1)) { + if (insn.Bits(24, 8) != 255 || (insn.words[2] & ~0x07800000U) != 0x007fe4ffU) + return fail(index, "native extended constant add mode"); + const u32 b = operand(1, false); + mov_imm(250, 1); + emit(0x5ca0000000000000ULL | predicate | 250 | (u64{250} << 8) | (u64{255} << 20) | + (u64{insn.Bits(87, 3)} << 39) | (u64{insn.Bits(90, 1)} << 42)); + emit(base(0x5c10000000000000ULL) | (u64{b} << 8) | (u64{250} << 20)); + break; + } + if (insn.Bits(64, 8) != 255 && !insn.Bits(73, 4) && insn.Bits(77, 7) == 0x7f && insn.Bits(84, 7) == 0x7f) { + const u32 b = operand(1, false); + emit(0x5c10000000000000ULL | predicate | 250 | (u64{insn.Bits(24, 8)} << 8) | + (u64{b} << 20) | (u64{insn.Bits(63, 1)} << 48) | (u64{insn.Bits(72, 1)} << 49)); + emit(base(0x5c10000000000000ULL) | (u64{250} << 8) | (u64{insn.Bits(64, 8)} << 20)); + break; + } + if (insn.Bits(64, 8) != 255 || insn.Bits(73, 4) || insn.Bits(77, 7) != 0x7f || insn.Bits(84, 7) != 0x7f) return fail(index, "native integer carry/third source"); + if (opcode == 0xa10 || opcode == 0xc10) { + if (opcode == 0xc10 && (insn.words[1] & ~0x8000003fU)) return fail(index, "native uniform add operand"); + const u32 b = operand(1, false); + emit(base(0x5c10000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | (u64{b} << 20) | + (u64{insn.Bits(63, 1)} << 48) | (u64{insn.Bits(72, 1)} << 49)); + } else if (opcode == 0x210) { + if (insn.words[1] & ~0x800000ffU) return fail(index, "native integer add register modifiers"); + emit(base(0x5c10000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{insn.Bits(32, 8)} << 20) | (u64{insn.Bits(63, 1)} << 48) | (u64{insn.Bits(72, 1)} << 49)); + } else { + emit(base(0x1c00000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | (u64{insn.words[1]} << 20) | (u64{insn.Bits(72, 1)} << 56)); + } + break; + } + case 0x949: case 0x958: { // Captured BRX / BRXU with absolute native-PC table values + if (insn.Bits(12, 4) != 7 || (insn.words[2] & ~0x3ffffU) != (opcode == 0x958 ? 0x0b800000U : 0x03800000U)) + return fail(index, "native computed branch mode"); + const u64 raw = insn.words[1] | (u64{insn.words[2] & 0x3ffff} << 32); + const s64 relative = static_cast(raw << 14) >> 14; + if (s64{index} * 16 + 16 + relative != 0) + return fail(index, "native computed branch relative base"); + const s64 offset = -s64{address_of(static_cast(result.instructions.size()))} - 8; + if (offset < -0x800000) return fail(index, "native computed branch range"); + const u32 source = opcode == 0x958 ? uniform_register(insn.Bits(24, 8)) : insn.Bits(24, 8); + if (source == ~0U) return fail(index, "native uniform computed branch source"); + emit(0xe25000000000000fULL | predicate | (u64{source} << 8) | + ((static_cast(offset) & 0xffffff) << 20)); + break; + } + case 0x944: { // Verified CALL.REL.NOINC with R55 holding the return PC + // A predicated NOINC forward transfer can terminate a loop at its + // reconvergence point without a return-link setup or matching RET. + if (native_predicate != 7 && insn.words[1] == 16 && insn.words[2] == 0x03c00000 && + index + 2 < code.size() && code[index + 1].Bits(0, 12) == 0x947) { + relocations.push_back({static_cast(result.instructions.size()), index + 2}); + emit(0xe24000000000000fULL | predicate); + break; + } + if (main_exited || index == 0 || insn.words[0] != 0x7944 || + (insn.words[2] & ~0x3ffffU) != 0x03c00000 || + (code[index - 1].words[0] & ~0x00ff0000U) != 0x00007802 || + code[index - 1].Bits(16, 8) >= 250 || + code[index - 1].words[1] != (index + 1) * 16 || + code[index - 1].words[2] != 0xf00) + return fail(index, "native call linkage mode"); + const u64 raw = insn.words[1] | (u64{insn.words[2] & 0x3ffff} << 32); + const s64 relative = static_cast(raw << 14) >> 14; + const s64 target = s64{index} * 16 + 16 + relative; + if (target < 0 || target % 16) return fail(index, "native call destination alignment"); + const u32 link = code[index - 1].Bits(16, 8); + for (size_t call = 0; call < call_targets.size(); ++call) { + if (call_targets[call] == target / 16 && call_links[call] != link) + return fail(index, "native inconsistent call linkage"); + } + call_targets.push_back(static_cast(target / 16)); + call_links.push_back(link); + relocations.push_back({static_cast(result.instructions.size()), target / 16}); + emit(0xe26000000007000fULL); + break; + } + case 0x950: { // Verified RET.REL.NODEC epilogue restoring the call link + const u32 return_reg = insn.Bits(24, 8); + if (!main_exited || index < 2 || !active_link || + (insn.words[0] & 0x00ffffffU) != 0x00007950 || return_reg >= 249 || + (insn.words[2] & ~0x3ffffU) != 0x03c00000 || + code[index - 2].words[0] != (0x00007202U | (return_reg << 16)) || + code[index - 2].words[1] != *active_link || code[index - 2].words[2] != 0xf00 || + code[index - 1].words[0] != (0x00007802U | ((return_reg + 1) << 16)) || + code[index - 1].words[1] || code[index - 1].words[2] != 0xf00 || !barrier_tokens.empty()) + return fail(index, "native return linkage mode"); + const u64 raw = insn.words[1] | (u64{insn.words[2] & 0x3ffff} << 32); + const s64 relative = static_cast(raw << 14) >> 14; + if (s64{index} * 16 + 16 + relative != 0) + return fail(index, "native return relative base"); + for (const auto& relocation : relocations) { + if (relocation.target == index || relocation.target == index - 1) + return fail(index, "native branch bypasses return linkage"); + } + emit(0xe32000000007000fULL); + bool complete = true; + for (u32 target : call_targets) complete &= target <= index; + if (complete) return finish(index); + break; + } + case 0x945: case 0x941: case 0x942: case 0x948: // (+ WARPSYNC) + // BSSY / BSYNC / BREAK only manage warp reconvergence: BSYNC waits for the other + // threads and falls through, BREAK drops the thread from that wait. Eden runs threads + // independently, so all three are no-ops (a SYNC would wrongly jump, and a branch + // out of a lowered SSY region would leave its token on Eden's flow stack). + break; + case 0x947: { // BRA + u64 branch_predicate = predicate; + if (opcode == 0x947 && insn.Bits(87, 4) != 7) { + if (insn.Bits(12, 4) != 7) { + relocations.push_back({static_cast(result.instructions.size()), index + 1}); + emit(0xe24000000000000fULL | (u64{insn.Bits(12, 4) ^ 8U} << 16)); + } + branch_predicate = u64{insn.Bits(87, 4)} << 16; + } else if (insn.Bits(87, 4) != 7) { + return fail(index, "native extra branch condition"); + } + const u64 raw = insn.words[1] | (u64{insn.words[2] & 0x3ffff} << 32); + const s64 relative = static_cast(raw << 14) >> 14; + const s64 target = s64{index} * 16 + 16 + relative; + if (target < 0 || target % 16) return fail(index, "native branch destination alignment"); + relocations.push_back({static_cast(result.instructions.size()), target / 16}); + emit(0xe24000000000000fULL | branch_predicate); + break; + } + case 0x94d: { // EXIT + if (insn.Bits(84, 2) || insn.Bits(87, 4) != 7) { + return fail(index, "native conditional/special EXIT"); + } + result.instructions.push_back(0xe30000000000000fULL | predicate); + if (insn.Bits(12, 4) != 7) break; + if (!barrier_tokens.empty()) return fail(index, "native unclosed reconvergence scope"); + if (call_targets.empty()) { + const bool pending_branch = std::any_of(relocations.begin(), relocations.end(), + [&](const Relocation& relocation) { return relocation.target > index; }); + if (!pending_branch) return finish(index); + break; + } + if (main_exited) return fail(index, "native library exits without returning"); + for (u32 target : call_targets) { + if (target <= index) return fail(index, "native call target overlaps main program"); + } + main_exited = true; + break; + } + default: + return fail(index, "unsupported native opcode"); + } + if (!operand_error.empty()) return fail(index, operand_error); + if (guarded_uniform) restore_predicate(); + } + return fail(static_cast(code.size()), main_exited ? "native program lacks captured library return" : "native program has no unconditional EXIT"); +} + +} // namespace Shader::Ampere diff --git a/src/shader_recompiler/frontend/ampere/program_address_map.h b/src/shader_recompiler/frontend/ampere/program_address_map.h new file mode 100644 index 0000000000..bee32a821c --- /dev/null +++ b/src/shader_recompiler/frontend/ampere/program_address_map.h @@ -0,0 +1,29 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later +#pragma once +#include +#include +#include +#include "common/common_types.h" + +namespace Shader::Ampere { +constexpr u64 NativeAddressMapMagic = 0x50414d4e4544454eULL; +inline bool HasNativeAddressMap(std::span code) { + return code.size() >= 2 && code.back() == NativeAddressMapMagic; +} +inline void AppendNativeAddressMap(std::vector& code, std::span offsets) { + if (offsets.empty()) return; + for (u32 index : offsets) code.push_back(0x50U + index * 8 + ((index + 1) / 3) * 8); + code.push_back(offsets.size()); + code.push_back(NativeAddressMapMagic); +} +inline std::optional NativeProgramAddress(std::span code, u32 address) { + if (!HasNativeAddressMap(code) || address % 16) return {}; + const u64 count = code[code.size() - 2]; + if (!count || count > 65536 || count + 2 > code.size() || address / 16 >= count) return {}; + const size_t start = code.size() - static_cast(count) - 2; + const u64 target = code[start + address / 16]; + if (target < 0x50 || target >= start * 8 || target % 8 || target % 32 == 0) return {}; + return static_cast(target); +} +} // namespace Shader::Ampere diff --git a/src/shader_recompiler/frontend/ampere/program_capture.h b/src/shader_recompiler/frontend/ampere/program_capture.h new file mode 100644 index 0000000000..ea7dfb7ef6 --- /dev/null +++ b/src/shader_recompiler/frontend/ampere/program_capture.h @@ -0,0 +1,43 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later +#pragma once + +#include "shader_recompiler/frontend/ampere/lower_to_maxwell.h" + +namespace Shader::Ampere { +class ProgramCaptureState { +public: + enum class Result { Continue, Complete, InvalidTarget }; + static constexpr u32 MaximumInstructions = 65536; + + Result Observe(const Instruction& instruction, u32 index) { + const u32 opcode = instruction.Bits(0, 12); + if (opcode == 0x944 || opcode == 0x947 || opcode == 0x945) { + const u64 raw = instruction.words[1] | (u64{instruction.words[2] & 0x3ffff} << 32); + const s64 relative = opcode == 0x945 ? static_cast(instruction.words[1]) + : static_cast(raw << 14) >> 14; + const s64 target = s64{index} * 16 + 16 + relative; + if (target < 0 || target % 16 || target >= MaximumInstructions * 16) + return Result::InvalidTarget; + if (opcode == 0x944) { + last_call_target = std::max(last_call_target, static_cast(target / 16)); + } else { + last_forward_target = std::max(last_forward_target, static_cast(target / 16)); + } + } + if (opcode == 0x94d && instruction.Bits(12, 4) == 7 && index >= last_forward_target) { + if (!last_call_target) return Result::Complete; + main_exited = true; + } + if (opcode == 0x950 && instruction.Bits(12, 4) == 7 && main_exited && + index >= last_call_target && index >= last_forward_target) + return Result::Complete; + return Result::Continue; + } + +private: + u32 last_call_target{}; + u32 last_forward_target{}; + bool main_exited{}; +}; +} // namespace Shader::Ampere diff --git a/src/shader_recompiler/frontend/ampere/program_override.h b/src/shader_recompiler/frontend/ampere/program_override.h new file mode 100644 index 0000000000..0943028b0c --- /dev/null +++ b/src/shader_recompiler/frontend/ampere/program_override.h @@ -0,0 +1,85 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later +#pragma once + +#include +#include "shader_recompiler/frontend/ampere/lower_to_maxwell.h" + +namespace Shader::Ampere { + +inline u64 NativeProgramKey(std::span source) { + u64 hash = 14695981039346656037ULL; + for (const auto& instruction : source) { + for (u32 word : instruction.words) { + for (u32 byte = 0; byte < 4; ++byte) hash = (hash ^ ((word >> (byte * 8)) & 255)) * 1099511628211ULL; + } + } + return hash; +} + +// Private iteration format: version, lengths, texture bank, exact original +// instructions, then the newly generated canonical instructions. Source bytes +// are compared directly; the filename hash is only an index. +inline std::vector EncodeProgramOverride(std::span source, + const LoweredProgram& program) { + if (!program || program.source_size != source.size_bytes() || program.source_instruction_offsets.size() != source.size()) return {}; + const std::array header{0x4e454445, 0x3230564f, + static_cast(source.size_bytes()), static_cast(program.instructions.size()), + program.required_texture_buffer, static_cast(program.source_instruction_offsets.size())}; + std::vector bytes(sizeof(header) + source.size_bytes() + program.instructions.size() * 8 + program.source_instruction_offsets.size() * 4); + std::memcpy(bytes.data(), header.data(), sizeof(header)); + std::memcpy(bytes.data() + sizeof(header), source.data(), source.size_bytes()); + std::memcpy(bytes.data() + sizeof(header) + source.size_bytes(), program.instructions.data(), program.instructions.size() * 8); + std::memcpy(bytes.data() + sizeof(header) + source.size_bytes() + program.instructions.size() * 8, + program.source_instruction_offsets.data(), program.source_instruction_offsets.size() * 4); + return bytes; +} + +inline u32 ProgramOverrideSourceSize(std::span bytes) { + std::array header{}; + if (bytes.size() < sizeof(header)) return 0; + std::memcpy(header.data(), bytes.data(), sizeof(header)); + const bool mapped = header[1] == 0x3230564f; + if (header[0] != 0x4e454445 || (!mapped && header[1] != 0x3130564f) || + (mapped ? header[5] != header[2] / 16 : header[5] != 0) || + header[2] == 0 || header[2] % sizeof(Instruction) || header[2] > 1048576 || + header[3] == 0 || header[3] > 1048576 || (header[4] != ~0U && header[4] >= 18)) return 0; + if (bytes.size() != sizeof(header) + size_t{header[2]} + size_t{header[3]} * 8 + size_t{header[5]} * 4) return 0; + return header[2]; +} + +inline LoweredProgram DecodeProgramOverride(std::span bytes, + std::span source) { + LoweredProgram result; + const auto fail = [&result](std::string_view message) { + result.instructions.clear(); + result.source_instruction_offsets.clear(); + result.error = message; + return result; + }; + std::array header{}; + if (!ProgramOverrideSourceSize(bytes)) return fail("invalid native override header or length"); + std::memcpy(header.data(), bytes.data(), sizeof(header)); + if (header[2] != source.size_bytes()) return fail("native override source length mismatch"); + if (std::memcmp(bytes.data() + sizeof(header), source.data(), source.size_bytes())) + return fail("native override source mismatch"); + result.instructions.resize(header[3]); + std::memcpy(result.instructions.data(), bytes.data() + sizeof(header) + source.size_bytes(), header[3] * 8); + if (header[5]) { + result.source_instruction_offsets.resize(header[5]); + std::memcpy(result.source_instruction_offsets.data(), + bytes.data() + sizeof(header) + source.size_bytes() + size_t{header[3]} * 8, + size_t{header[5]} * 4); + u32 previous = 0; + for (u32 offset : result.source_instruction_offsets) { + if (offset >= result.instructions.size() || offset < previous) + return fail("native override address map mismatch"); + previous = offset; + } + } + result.source_size = header[2]; + result.required_texture_buffer = header[4]; + return result; +} + +} // namespace Shader::Ampere diff --git a/src/shader_recompiler/frontend/maxwell/control_flow.cpp b/src/shader_recompiler/frontend/maxwell/control_flow.cpp index 7dabfbf16e..9fb30df664 100644 --- a/src/shader_recompiler/frontend/maxwell/control_flow.cpp +++ b/src/shader_recompiler/frontend/maxwell/control_flow.cpp @@ -457,7 +457,7 @@ CFG::AnalysisState CFG::AnalyzeBRX(Block* block, Location pc, Instruction inst, block->indirect_branches.reserve(targets.size()); for (const u32 target : targets) { - Block* const branch{AddLabel(block, block->stack, target, function_id)}; + Block* const branch{AddLabel(block, block->stack, env.TranslateProgramAddress(target), function_id)}; block->indirect_branches.push_back({ .block = branch, .address = target, diff --git a/src/video_core/CMakeLists.txt b/src/video_core/CMakeLists.txt index 312a6a09db..d17deacacd 100644 --- a/src/video_core/CMakeLists.txt +++ b/src/video_core/CMakeLists.txt @@ -90,6 +90,7 @@ add_library(video_core STATIC host1x/vic.h macro.cpp macro.h + macro_turing.h fence_manager.h gpu.cpp gpu.h @@ -99,6 +100,8 @@ add_library(video_core STATIC invalidation_accumulator.h memory_manager.cpp memory_manager.h + native_qmd.h + native_query_trace.h nx2.h present.h pte_kind.h diff --git a/src/video_core/buffer_cache/buffer_cache.h b/src/video_core/buffer_cache/buffer_cache.h index f30c85911f..126e73c9b6 100644 --- a/src/video_core/buffer_cache/buffer_cache.h +++ b/src/video_core/buffer_cache/buffer_cache.h @@ -9,6 +9,7 @@ #include #include #include +#include #include "common/range_sets.inc" #include "video_core/buffer_cache/buffer_cache_base.h" @@ -1372,7 +1373,12 @@ void BufferCache

::UpdateVertexBuffer(u32 index) { const auto& array = maxwell3d->regs.vertex_streams[index]; const auto& limit = maxwell3d->regs.vertex_stream_limits[index]; const GPUVAddr gpu_addr_begin = array.Address(); - const GPUVAddr gpu_addr_end = limit.Address() + 1; + const bool native = maxwell3d->GetNativeMacroEngine().native_turing; + const u64 native_size = native ? maxwell3d->regs.NativeVertexStreamSize(index) : 0; + if (native && array.enable != 0 && native_size > (std::numeric_limits::max)()) { + throw std::runtime_error("Native vertex stream exceeds supported 32-bit buffer size"); + } + const GPUVAddr gpu_addr_end = native ? gpu_addr_begin + native_size : limit.Address() + 1; const std::optional device_addr = gpu_memory->GpuToCpuAddress(gpu_addr_begin); const u32 address_size = static_cast(gpu_addr_end - gpu_addr_begin); u32 size = address_size; // TODO: Analyze stride and number of vertices diff --git a/src/video_core/control/channel_state.cpp b/src/video_core/control/channel_state.cpp index a1a8110510..57d7dc8e57 100644 --- a/src/video_core/control/channel_state.cpp +++ b/src/video_core/control/channel_state.cpp @@ -13,7 +13,9 @@ #include "video_core/engines/maxwell_3d.h" #include "video_core/engines/maxwell_dma.h" #include "video_core/engines/puller.h" +#include "video_core/macro.h" #include "video_core/memory_manager.h" +#include "video_core/nx2.h" namespace Tegra::Control { namespace { @@ -25,6 +27,17 @@ constexpr u32 NvkM2mfSubchannel = 2; constexpr u32 Nvk2DSubchannel = 3; constexpr u32 NvkCopySubchannel = 4; +// Switch 2 channels of a process share one TSG, whose graphics context holds the macro RAM. +std::shared_ptr Nx2MacroContext() { + static const auto context = [] { + auto created = std::make_shared(); + created->engine.native_turing = true; + created->compute_engine.native_turing = true; + return created; + }(); + return context; +} + void BindNvkDefaultSubchannels(ChannelState::Payload& payload) { auto& dma_pusher = payload.dma_pusher; dma_pusher.BindSubchannel(&payload.maxwell_3d, Nvk3DSubchannel, Engines::EngineTypes::Maxwell3D); @@ -60,6 +73,16 @@ void ChannelState::Init(Core::System& system, u64 program_id_) { program_id = program_id_; payload.emplace(system, *memory_manager, *this); BindNvkDefaultSubchannels(*payload); + if (g_nx2_ampere) { + // Switch 2 command streams use Ampere classes with Turing+ (MME64) macros, and its + // block-linear surfaces the console's GOB swizzle. + memory_manager->EnableNativeTiling(); + const auto macros = Nx2MacroContext(); + payload->maxwell_3d.SetNativeMacroContext(macros); + payload->maxwell_3d.SetNativeTuringMacroMode(true); + payload->kepler_compute.SetNativeMacroMode(¯os->engine, macros->positions, + macros.get()); + } initialized = true; } diff --git a/src/video_core/dirty_flags.cpp b/src/video_core/dirty_flags.cpp index c2ecc12f56..e6856aaca3 100644 --- a/src/video_core/dirty_flags.cpp +++ b/src/video_core/dirty_flags.cpp @@ -22,6 +22,7 @@ void SetupDirtyVertexBuffers(Maxwell3D::DirtyState::Tables& tables) { OFF(vertex_stream_limits) + i * NUM(vertex_stream_limits[0]); FillBlock(tables, array_offset, num_array, VertexBuffer0 + i, VertexBuffers); + FillBlock(tables, 0x600 / sizeof(u32) + i * 2, 2, VertexBuffer0 + i, VertexBuffers); FillBlock(tables, limit_offset, NUM(vertex_stream_limits), VertexBuffer0 + i, VertexBuffers); } diff --git a/src/video_core/engines/engine_upload.cpp b/src/video_core/engines/engine_upload.cpp index ba21eec4c4..0ffdae8036 100644 --- a/src/video_core/engines/engine_upload.cpp +++ b/src/video_core/engines/engine_upload.cpp @@ -27,6 +27,7 @@ void State::BindRasterizer(VideoCore::RasterizerInterface* rasterizer_) { void State::ProcessExec(const bool is_linear_) { write_offset = 0; + flushed_offset = 0; copy_size = regs.line_length_in * regs.line_count; inner_buffer.resize_destructive(copy_size); is_linear = is_linear_; @@ -39,10 +40,29 @@ void State::ProcessData(const u32 data, const bool is_last_call) { if (!is_last_call) { return; } + if (native_partial && is_linear && regs.line_count == 1) { + const u32 count = write_offset - flushed_offset; + if (count) { + rasterizer->AccelerateInlineToMemory(regs.dest.Address() + flushed_offset, count, + std::span{inner_buffer.data() + flushed_offset, count}); + flushed_offset = write_offset; + } + return; + } ProcessData(inner_buffer); } void State::ProcessData(const u32* data, size_t num_data) { + if (native_partial && is_linear && regs.line_count == 1) { + const auto count = std::min(num_data * sizeof(u32), copy_size - write_offset); + if (count) { + rasterizer->AccelerateInlineToMemory(regs.dest.Address() + write_offset, count, + std::span{reinterpret_cast(data), count}); + write_offset += static_cast(count); + flushed_offset = write_offset; + } + return; + } std::span read_buffer(reinterpret_cast(data), num_data * sizeof(u32)); ProcessData(read_buffer); } diff --git a/src/video_core/engines/engine_upload.h b/src/video_core/engines/engine_upload.h index 21bf8aeb46..55a82f938b 100644 --- a/src/video_core/engines/engine_upload.h +++ b/src/video_core/engines/engine_upload.h @@ -62,6 +62,7 @@ public: explicit State(MemoryManager& memory_manager_, Registers& regs_); ~State(); + void SetNativePartialTransfers(bool enabled) { native_partial = enabled; } void ProcessExec(bool is_linear_); void ProcessData(u32 data, bool is_last_call); void ProcessData(const u32* data, size_t num_data); @@ -80,6 +81,8 @@ public: private: void ProcessData(std::span read_buffer); + bool native_partial{}; + u32 flushed_offset{}; u32 write_offset = 0; u32 copy_size = 0; Common::ScratchBuffer inner_buffer; diff --git a/src/video_core/engines/fermi_2d.cpp b/src/video_core/engines/fermi_2d.cpp index aeff5e156f..a9b3355476 100644 --- a/src/video_core/engines/fermi_2d.cpp +++ b/src/video_core/engines/fermi_2d.cpp @@ -4,6 +4,7 @@ // SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include #include "common/assert.h" #include "common/logging.h" #include "video_core/engines/fermi_2d.h" @@ -59,6 +60,14 @@ void Fermi2D::ConsumeSinkImpl(Core::System& system) { } void Fermi2D::Blit() { + if (static const bool trace = std::getenv("EDEN_NX2_RT_TRACE") != nullptr; trace) { + static u32 traced{}; + if (traced++ < 64) { + LOG_INFO(HW_GPU, "nx2 2D blit {:#x} {}x{} -> {:#x} {}x{}", regs.src.Address(), + regs.src.width, regs.src.height, regs.dst.Address(), regs.dst.width, + regs.dst.height); + } + } LOG_DEBUG(HW_GPU, "called. source address={:#x}, destination address={:#x}", regs.src.Address(), regs.dst.Address()); UNIMPLEMENTED_IF_MSG(regs.operation != Operation::SrcCopy, "Operation is not copy"); diff --git a/src/video_core/engines/kepler_compute.cpp b/src/video_core/engines/kepler_compute.cpp index 166bbd6d53..7710ea7dd2 100644 --- a/src/video_core/engines/kepler_compute.cpp +++ b/src/video_core/engines/kepler_compute.cpp @@ -5,12 +5,16 @@ // SPDX-License-Identifier: GPL-2.0-or-later #include +#include +#include #include "common/assert.h" #include "common/logging.h" +#include "shader_recompiler/exception.h" #include "core/core.h" #include "video_core/engines/kepler_compute.h" #include "video_core/engines/maxwell_3d.h" #include "video_core/memory_manager.h" +#include "video_core/native_query_trace.h" #include "video_core/rasterizer_interface.h" #include "video_core/textures/decoders.h" @@ -42,15 +46,175 @@ void KeplerCompute::ConsumeSinkImpl(Core::System& system) { method_sink.clear(); } +bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 argument, bool is_last_call) { + if (native_context) { + if (method == 0x45 || method == 0x47) native_compute_uploaded = true; + native_macro = native_compute_uploaded ? &native_context->compute_engine : &native_context->engine; + native_macro_positions = native_compute_uploaded ? std::span{native_context->compute_positions} : std::span{native_context->positions}; + } + if (method >= native_registers.size()) { + LOG_ERROR(HW_GPU, "Native compute method out of range: {:#x}", method); + return true; + } + if (method >= 0xe00) { + if (!native_executing_macro) { + if (method & 1) { LOG_ERROR(HW_GPU, "Native compute macro starts on argument method"); return true; } + native_executing_macro = method; + } + native_parameters.push_back(argument); + if (!is_last_call) return true; + const u32 entry = ((*native_executing_macro - 0xe00) / 2) % native_macro_positions.size(); + native_executing_macro.reset(); + const auto result = native_macro->native_interpreter.Execute(native_macro->native_code, + native_macro_positions[entry], native_parameters, + [&](u32 address) { return address < native_registers.size() ? native_registers[address] : 0U; }, + [&](u32 address, u32 value) { + static u32 diagnostic_writes{}; + if (diagnostic_writes++ < 4096) LOG_DEBUG(HW_GPU, "Native compute MME emit method={:#x}, value={:#x}", address, value); + if (address >= 0xe00) return false; + CallMethod(system, address, value, true); + return true; + }); + if (!result.success) { + native_macro->native_execution_failed = true; + if (std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS")) { + std::ofstream dump("native-compute-mme.bin", std::ios::binary); + dump.write(reinterpret_cast(native_macro->native_code.data()), sizeof(native_macro->native_code)); + } + LOG_ERROR(HW_GPU, "Native MME failed: compute start={:#x}, pc={:#x}, consumed={}/{}, error={}, parameters={:08x}", + native_macro_positions[entry], result.pc, result.consumed, native_parameters.size(), result.error, + fmt::join(native_parameters, ",")); + } + native_parameters.clear(); + return true; + } + native_registers[method] = argument; + if (method == 0x5a || (method >= 0xc6 && method <= 0x108)) { + VideoCommon::TraceNativeQuery(12, memory_manager.GetID(), method, 0, 0, argument, 0); + } + if (method == 0x107) { + const GPUVAddr address = ((GPUVAddr{native_registers[0xc6]} << 32) | + native_registers[0xc7]) << 8; + const GPUVAddr program = (GPUVAddr{native_registers[0xc8 + 49] & 0x1ffff} << 32) | + native_registers[0xc8 + 48]; + const u32 control = native_registers[0xa7]; + LOG_WARNING(HW_GPU, "Native inline QMD: address={:#x}, program={:#x}, control={:#x}, grid={}/{}/{}", + address, program, control, native_registers[0xc8 + 12], + native_registers[0xc8 + 13] & 0xffff, native_registers[0xc8 + 14] & 0xffff); + if (std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS")) { + std::ofstream dump("native-inline-qmd.bin", std::ios::binary); + dump.write(reinterpret_cast(native_registers.data() + 0xc8), 256); + std::ofstream registers("native-inline-compute-registers.bin", std::ios::binary); + registers.write(reinterpret_cast(native_registers.data()), sizeof(native_registers)); + if (memory_manager.IsFullyMappedRange(program, 0x10000)) { + std::vector code(0x10000); + memory_manager.ReadBlockUnsafe(program, code.data(), code.size()); + std::ofstream shader("native-inline-compute-shader.bin", std::ios::binary); + shader.write(reinterpret_cast(code.data()), code.size()); + } + } + if (!memory_manager.IsFullyMappedRange(address, 256)) { + native_macro->native_execution_failed = true; + LOG_ERROR(HW_GPU, "Unmapped native inline QMD destination"); + return true; + } + rasterizer->AccelerateInlineToMemory(address, 256, + std::span{reinterpret_cast(native_registers.data() + 0xc8), 256}); + if ((control & (1U << 24)) || !(native_registers[0xcc] & (1U << 12))) { + const u32 previous = native_registers[0xad]; + native_registers[0xad] = static_cast(address >> 8); + ProcessNativeMethod(system, 0xb0, 3, true); + native_registers[0xad] = previous; + } + return true; + } + switch (method) { + case 0x45: native_macro->ClearCode(argument); return true; + case 0x46: native_macro->AddCode(0, argument); return true; + case 0x47: native_bind_pointer = argument; return true; + case 0x48: + if (native_bind_pointer < native_macro_positions.size()) native_macro_positions[native_bind_pointer++] = argument; + else LOG_ERROR(HW_GPU, "Native compute macro bind out of bounds"); + return true; + case 0xae: case 0xaf: case 0xb0: { + const GPUVAddr qmd_address = GPUVAddr{native_registers[0xad]} << 8; + VideoCommon::TraceNativeQuery(9, memory_manager.GetID(), qmd_address, + memory_manager.GpuToCpuAddress(qmd_address).value_or(0), method, argument, 0); + if (method == 0xae || (method == 0xaf && !(argument & 2)) || + (method == 0xb0 && ((argument & 15) == 0 || (argument & 15) == 1))) return true; + if ((argument & 15) == 12 && std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS")) { + static u32 captured{}; + if (captured++ == 0 && memory_manager.IsFullyMappedRange(qmd_address, 256)) { + std::array descriptor{}; + memory_manager.ReadBlockUnsafe(qmd_address, descriptor.data(), sizeof(descriptor)); + std::ofstream dump("native-increment-qmd.bin", std::ios::binary); + dump.write(reinterpret_cast(descriptor.data()), sizeof(descriptor)); + LOG_WARNING(HW_GPU, "Native increment dependency descriptor: address={:#x}, words={:08x}", + qmd_address, fmt::join(descriptor, ",")); + } + } + const auto error = native_qmd.Signal(qmd_address, argument, + [&](u64 address, NativeQmdScheduler::Descriptor& descriptor) { + if (!memory_manager.IsFullyMappedRange(address, sizeof(descriptor))) return false; + memory_manager.ReadBlockUnsafe(address, descriptor.data(), sizeof(descriptor)); + for (u32 slot = 0; slot < 2; ++slot) { + const u32 flags = descriptor[25 + slot * 4]; + if (!(flags & (1U << 23))) continue; + const GPUVAddr destination = (GPUVAddr{flags & 0xff} << 32) | descriptor[24 + slot * 4]; + VideoCommon::TraceNativeQuery(7 + slot, memory_manager.GetID(), destination, + memory_manager.GpuToCpuAddress(destination).value_or(0), + descriptor[4], descriptor[26 + slot * 4], flags); + } + return true; + }, + [&](u64 address, u32 payload, u32 flags) { + if (!memory_manager.IsFullyMappedRange(address, sizeof(payload))) return false; + if (flags & (1U << 26)) { + rasterizer->SignalFence([this, address, payload] { + const u32 value = memory_manager.Read(address) + payload; + memory_manager.Write(address, value); + VideoCommon::TraceNativeQuery(1, memory_manager.GetID(), address, + memory_manager.GpuToCpuAddress(address).value_or(0), 0, value, 1U << 26); + }); + } else { + rasterizer->Query(address, VideoCommon::QueryType::Payload, + VideoCommon::QueryPropertiesFlags::IsAFence, payload, 0); + } + return true; + }, [&](const auto& descriptor) { return DispatchNativeCompute(descriptor); }); + if (error.empty()) { + LOG_DEBUG(HW_GPU, "Native QMD control action={:#x}, QMD={:#x}", argument, qmd_address); + return true; + } + native_macro->native_execution_failed = true; + LOG_ERROR(HW_GPU, "Native QMD failed: {}", error); + std::array qmd{}; + if (memory_manager.IsFullyMappedRange(qmd_address, sizeof(qmd))) { + memory_manager.ReadBlockUnsafe(qmd_address, qmd.data(), sizeof(qmd)); + } + if (std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS") && memory_manager.IsFullyMappedRange(qmd_address, sizeof(qmd))) { + std::ofstream dump("native-qmd.bin", std::ios::binary); + dump.write(reinterpret_cast(qmd.data()), sizeof(qmd)); + std::ofstream macros("native-dispatch-mme.bin", std::ios::binary); + macros.write(reinterpret_cast(native_macro->native_code.data()), sizeof(native_macro->native_code)); + std::ofstream registers("native-dispatch-registers.bin", std::ios::binary); + registers.write(reinterpret_cast(native_registers.data()), sizeof(native_registers)); + } + LOG_ERROR(HW_GPU, "Native compute dispatch unsupported: method={:#x}, value={:#x}, QMD={:#x}, words={:08x}", method, argument, + qmd_address, fmt::join(qmd, ",")); + return true; + } + default: + // The inline upload register block is shared with the older class. + if (method < Regs::NUM_REGS) regs.reg_array[method] = argument; + return method != KEPLER_COMPUTE_REG_INDEX(exec_upload) && method != KEPLER_COMPUTE_REG_INDEX(data_upload); + } +} + void KeplerCompute::CallMethod(Core::System& system, u32 method, u32 method_argument, bool is_last_call) { + if (native_mode && ProcessNativeMethod(system, method, method_argument, is_last_call)) return; if (method >= Regs::NUM_REGS) { - // Ampere compute (Switch 2) uses methods past Kepler's register file; writing them would - // run off reg_array, so drop them. - static bool reported = false; - if (!reported) { - reported = true; - LOG_WARNING(HW_GPU, "Ignoring KeplerCompute method {:#x} past the register file", method); - } + LOG_ERROR(HW_GPU, "Compute register out of range: {:#x}", method); return; } @@ -96,6 +260,10 @@ void KeplerCompute::CallMethod(Core::System& system, u32 method, u32 method_argu } void KeplerCompute::CallMultiMethod(Core::System& system, u32 method, const u32* base_start, u32 amount, u32 methods_pending) { + if (native_mode) { + for (u32 i = 0; i < amount; ++i) CallMethod(system, method, base_start[i], methods_pending - i <= 1); + return; + } switch (method) { case KEPLER_COMPUTE_REG_INDEX(data_upload): upload_address = current_dma_segment; @@ -111,6 +279,42 @@ void KeplerCompute::CallMultiMethod(Core::System& system, u32 method, const u32* } } +bool KeplerCompute::DispatchNativeCompute(const NativeQmdScheduler::Descriptor& qmd) { + launch_description = {}; + launch_description.grid_dim_x.Assign(qmd[12]); + launch_description.grid_dim_y.Assign(qmd[13] & 0xffff); + launch_description.grid_dim_z.Assign(qmd[14] & 0xffff); + launch_description.shared_alloc.Assign(qmd[17] & 0x3ffff); + launch_description.block_dim_x.Assign(qmd[18] >> 16); + launch_description.block_dim_y.Assign(qmd[19] & 0xffff); + launch_description.block_dim_z.Assign(qmd[19] >> 16); + launch_description.const_buffer_enable_mask.Assign(qmd[20] & 0xff); + launch_description.local_pos_alloc.Assign(qmd[23] & 0xffffff); + launch_description.local_crs_alloc.Assign(qmd[50] & 0xffffff); + regs.code_loc.address_low = qmd[48]; + regs.code_loc.address_high = qmd[49] & 0x1ffff; + for (u32 slot = 0; slot < NumConstBuffers; ++slot) { + const u32 high = qmd[33 + slot * 2]; + if (high & 0x1ff00) { + LOG_ERROR(HW_GPU, "Native compute constant address exceeds supported GPU width"); + return false; + } + auto& buffer = launch_description.const_buffer_config[slot]; + buffer.address_low = qmd[32 + slot * 2]; + buffer.address_high.Assign(high & 0xff); + buffer.size.Assign(((high >> 19) & 0x1fff) << 4); + } + if (!launch_description.block_dim_x || !launch_description.block_dim_y || !launch_description.block_dim_z) + return false; + try { + rasterizer->DispatchCompute(); + } catch (const Shader::Exception& error) { + LOG_ERROR(HW_GPU, "Native compute shader rejected: {}", error.what()); + return false; + } + return true; +} + void KeplerCompute::ProcessLaunch() { const GPUVAddr launch_desc_loc = regs.launch_desc_loc.Address(); memory_manager.ReadBlockUnsafe(launch_desc_loc, &launch_description, diff --git a/src/video_core/engines/kepler_compute.h b/src/video_core/engines/kepler_compute.h index 078d3d45be..99cf74d1e8 100644 --- a/src/video_core/engines/kepler_compute.h +++ b/src/video_core/engines/kepler_compute.h @@ -15,6 +15,8 @@ #include "common/common_types.h" #include "video_core/engines/engine_interface.h" #include "video_core/engines/engine_upload.h" +#include "video_core/macro.h" +#include "video_core/native_qmd.h" #include "video_core/textures/texture.h" namespace Core { @@ -51,6 +53,22 @@ public: /// Binds a rasterizer to this engine. void BindRasterizer(VideoCore::RasterizerInterface* rasterizer); + void SetNativeMacroMode(MacroEngine* shared_macro, std::span positions, NativeMacroContext* context = nullptr) { + if (native_context != context) native_compute_uploaded = false; + native_context = context; + native_mode = shared_macro != nullptr; + upload_state.SetNativePartialTransfers(native_mode); + native_macro = shared_macro; + native_macro_positions = positions; + if (native_mode) execution_mask.set(); + else { + execution_mask.reset(); + execution_mask[KEPLER_COMPUTE_REG_INDEX(exec_upload)] = true; + execution_mask[KEPLER_COMPUTE_REG_INDEX(data_upload)] = true; + execution_mask[KEPLER_COMPUTE_REG_INDEX(launch)] = true; + } + } + static constexpr std::size_t NumConstBuffers = 8; struct Regs { @@ -211,7 +229,22 @@ public: return indirect_compute; } + bool IsNativeMode() const { return native_mode; } + private: + bool DispatchNativeCompute(const NativeQmdScheduler::Descriptor& descriptor); + bool ProcessNativeMethod(Core::System& system, u32 method, u32 argument, bool is_last_call); + NativeQmdScheduler native_qmd; + bool native_mode{}; + NativeMacroContext* native_context{}; + bool native_compute_uploaded{}; + MacroEngine* native_macro{}; + std::array native_registers{}; + std::span native_macro_positions; + std::vector native_parameters; + std::optional native_executing_macro; + u32 native_bind_pointer{}; + void ProcessLaunch(); void ConsumeSinkImpl(Core::System& system) override; diff --git a/src/video_core/engines/maxwell_3d.cpp b/src/video_core/engines/maxwell_3d.cpp index 1a43133296..f86b8e9351 100644 --- a/src/video_core/engines/maxwell_3d.cpp +++ b/src/video_core/engines/maxwell_3d.cpp @@ -5,6 +5,7 @@ // SPDX-License-Identifier: GPL-2.0-or-later #include +#include #include #include "common/assert.h" #include "common/bit_util.h" @@ -12,7 +13,6 @@ #include "common/settings.h" #include "core/core.h" #include "core/core_timing.h" -#include "video_core/nx2.h" #include "video_core/dirty_flags.h" #include "video_core/engines/maxwell_3d.h" #include "video_core/engines/maxwell_3d.h" @@ -254,7 +254,10 @@ u32 Maxwell3D::GetMaxCurrentVertices() { const auto& limit = regs.vertex_stream_limits[index]; const GPUVAddr gpu_addr_begin = array.Address(); const GPUVAddr gpu_addr_end = limit.Address() + 1; - const u32 address_size = static_cast(gpu_addr_end - gpu_addr_begin); + const u64 stream_size = GetNativeMacroEngine().native_turing + ? regs.NativeVertexStreamSize(index) + : gpu_addr_end - gpu_addr_begin; + const u32 address_size = static_cast(std::min(stream_size, std::numeric_limits::max())); num_vertices = (std::max)( num_vertices, address_size / (std::max)(attribute.SizeInBytes(), array.stride.Value())); break; @@ -306,6 +309,9 @@ void Maxwell3D::ProcessDirtyRegisters(u32 method, u32 argument) { } void Maxwell3D::ProcessMethodCall(u32 method, u32 argument, u32 nonshadow_argument, bool is_last_call) { + if (GetNativeMacroEngine().native_turing && method == MAXWELL3D_REG_INDEX(clear_surface)) { + LOG_DEBUG(HW_GPU, "Native GPU surface clear: value={:#x}", argument); + } switch (method) { case MAXWELL3D_REG_INDEX(wait_for_idle): return rasterizer->WaitForIdle(); @@ -313,9 +319,9 @@ void Maxwell3D::ProcessMethodCall(u32 method, u32 argument, u32 nonshadow_argume shadow_state.shadow_ram_control = static_cast(nonshadow_argument); return; case MAXWELL3D_REG_INDEX(load_mme.instruction_ptr): - return macro_engine.ClearCode(regs.load_mme.instruction_ptr); + return GetNativeMacroEngine().ClearCode(regs.load_mme.instruction_ptr); case MAXWELL3D_REG_INDEX(load_mme.instruction): - return macro_engine.AddCode(regs.load_mme.instruction_ptr, argument); + return GetNativeMacroEngine().AddCode(regs.load_mme.instruction_ptr, argument); case MAXWELL3D_REG_INDEX(load_mme.start_address): return ProcessMacroBind(argument); case MAXWELL3D_REG_INDEX(falcon[4]): @@ -379,16 +385,10 @@ void Maxwell3D::CallMacroMethod(Core::System& system, u32 method, const std::vec // Lookup the macro offset const u32 entry = - ((method - MacroRegistersStart) >> 1) % static_cast(macro_positions.size()); - - // Switch 2 uploads Turing+ (MME64) macros, which Eden's Fermi-style macro engine can't run. - if (g_nx2_ampere) { - LOG_DEBUG(HW_GPU, "Switch 2 macro {} with {} parameters skipped", entry, parameters.size()); - return; - } + ((method - MacroRegistersStart) >> 1) % static_cast(GetNativeMacroPositions().size()); // Execute the current macro. - macro_engine.Execute(system, *this, macro_positions[entry], parameters); + GetNativeMacroEngine().Execute(system, *this, GetNativeMacroPositions()[entry], parameters); draw_manager.DrawDeferred(*this); } @@ -460,11 +460,11 @@ void Maxwell3D::CallMultiMethod(Core::System& system, u32 method, const u32* bas } void Maxwell3D::ProcessMacroUpload(u32 data) { - macro_engine.AddCode(regs.load_mme.instruction_ptr++, data); + GetNativeMacroEngine().AddCode(regs.load_mme.instruction_ptr++, data); } void Maxwell3D::ProcessMacroBind(u32 data) { - macro_positions[regs.load_mme.start_address_ptr++] = data; + GetNativeMacroPositions()[regs.load_mme.start_address_ptr++] = data; } void Maxwell3D::ProcessFirmwareCall4() { @@ -593,6 +593,9 @@ void Maxwell3D::ProcessCounterReset() { } void Maxwell3D::ProcessSyncPoint() { + if (GetNativeMacroEngine().native_turing) { + LOG_DEBUG(HW_GPU, "Native 3D syncpoint increment: id={}, raw={:#x}", regs.sync_info.sync_point.Value(), regs.reg_array[MAXWELL3D_REG_INDEX(sync_info)]); + } const u32 sync_point = regs.sync_info.sync_point.Value(); [[maybe_unused]] const u32 cache_flush = regs.sync_info.clean_l2.Value(); rasterizer->SignalSyncPoint(sync_point); diff --git a/src/video_core/engines/maxwell_3d.h b/src/video_core/engines/maxwell_3d.h index 0cab2ae986..8cbc4f4274 100644 --- a/src/video_core/engines/maxwell_3d.h +++ b/src/video_core/engines/maxwell_3d.h @@ -2605,6 +2605,12 @@ public: u32 stop_shader_mask; }; + // Native C597/C797 use byte counts at 0x600 rather than legacy end addresses. + u64 NativeVertexStreamSize(std::size_t index) const { + const std::size_t method = 0x600 / sizeof(u32) + index * 2; + return (u64{reg_array[method] & 0xffU} << 32) | reg_array[method + 1]; + } + // clang-format off union { struct { @@ -3129,6 +3135,17 @@ public: /// Reads a register value located at the input method address u32 GetRegisterValue(u32 method) const; + void SetNativeTuringMacroMode(bool enabled) { + GetNativeMacroEngine().native_turing = enabled; + upload_state.SetNativePartialTransfers(enabled); + } + void SetNativeMacroContext(std::shared_ptr context) { + native_macro_context = std::move(context); + } + MacroEngine& GetNativeMacroEngine() { return native_macro_context ? native_macro_context->engine : macro_engine; } + std::span GetNativeMacroPositions() { return native_macro_context ? std::span{native_macro_context->positions} : std::span{macro_positions}; } + + /// Write the value to the register identified by method. void CallMethod(Core::System& system, u32 method, u32 method_argument, bool is_last_call) override; @@ -3264,6 +3281,7 @@ private: std::vector macro_params; /// Interpreter for the macro codes uploaded to the GPU. + std::shared_ptr native_macro_context; MacroEngine macro_engine; Upload::State upload_state; diff --git a/src/video_core/engines/maxwell_dma.cpp b/src/video_core/engines/maxwell_dma.cpp index 5ab953e510..5b6efabca0 100644 --- a/src/video_core/engines/maxwell_dma.cpp +++ b/src/video_core/engines/maxwell_dma.cpp @@ -4,6 +4,7 @@ // SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include #include "common/algorithm.h" #include "common/assert.h" #include "common/logging.h" @@ -63,6 +64,14 @@ void MaxwellDMA::CallMultiMethod(Core::System& system, u32 method, const u32* ba void MaxwellDMA::Launch() { LOG_TRACE(Render_OpenGL, "DMA copy {:#x} -> {:#x}", static_cast(regs.offset_in), GPUVAddr(regs.offset_out)); + if (static const bool trace = std::getenv("EDEN_NX2_RT_TRACE") != nullptr; trace) { + static u32 traced{}; + if (traced++ < 64) { + LOG_INFO(HW_GPU, "nx2 DMA {:#x} -> {:#x} lines={} length={} launch={:#x}", + static_cast(regs.offset_in), GPUVAddr(regs.offset_out), + regs.line_count, regs.line_length_in, regs.reg_array[0xc0]); + } + } // TODO(Subv): Perform more research and implement all features of this engine. const LaunchDMA& launch = regs.launch_dma; @@ -223,7 +232,7 @@ void MaxwellDMA::CopyBlockLinearToPitch() { UnswizzleSubrect(tmp_write_buffer, tmp_read_buffer, bytes_per_pixel, width, height, depth, x_offset, src_params.origin.y, x_elements, regs.line_count, block_height, - block_depth, dst_operand.pitch); + block_depth, dst_operand.pitch, memory_manager.UsesNativeTiling()); } void MaxwellDMA::CopyPitchToBlockLinear() { @@ -288,7 +297,7 @@ void MaxwellDMA::CopyPitchToBlockLinear() { // If the input is linear and the output is tiled, swizzle the input and copy it over. SwizzleSubrect(tmp_write_buffer, tmp_read_buffer, bytes_per_pixel, width, height, depth, x_offset, dst_params.origin.y, x_elements, regs.line_count, block_height, - block_depth, regs.pitch_in); + block_depth, regs.pitch_in, memory_manager.UsesNativeTiling()); } void MaxwellDMA::CopyBlockLinearToBlockLinear() { @@ -341,11 +350,11 @@ void MaxwellDMA::CopyBlockLinearToBlockLinear() { UnswizzleSubrect(intermediate_buffer, tmp_read_buffer, bytes_per_pixel, src_width, src.height, src.depth, src_x_offset, src.origin.y, x_elements, regs.line_count, - src.block_size.height, src.block_size.depth, pitch); + src.block_size.height, src.block_size.depth, pitch, memory_manager.UsesNativeTiling()); SwizzleSubrect(tmp_write_buffer, intermediate_buffer, bytes_per_pixel, dst_width, dst.height, dst.depth, dst_x_offset, dst.origin.y, x_elements, regs.line_count, - dst.block_size.height, dst.block_size.depth, pitch); + dst.block_size.height, dst.block_size.depth, pitch, memory_manager.UsesNativeTiling()); } void MaxwellDMA::ReleaseSemaphore() { diff --git a/src/video_core/host_shaders/block_linear_unswizzle_2d.comp b/src/video_core/host_shaders/block_linear_unswizzle_2d.comp index 2a3446ca66..bbe73acd1a 100644 --- a/src/video_core/host_shaders/block_linear_unswizzle_2d.comp +++ b/src/video_core/host_shaders/block_linear_unswizzle_2d.comp @@ -39,6 +39,7 @@ UNIFORM(4) uint block_size; UNIFORM(5) uint x_shift; UNIFORM(6) uint block_height; UNIFORM(7) uint block_height_mask; +UNIFORM(8) uint native_tiling; END_PUSH_CONSTANTS #if HAS_EXTENDED_TYPES @@ -77,7 +78,12 @@ uint SwizzleTable(uint pos) { uint SwizzleOffset(uvec2 pos) { pos = pos & SWIZZLE_MASK; - return SwizzleTable(pos.y * 64 + pos.x); + uint offset = SwizzleTable(pos.y * 64 + pos.x); + if (native_tiling != 0) { + uint difference = ((offset >> 5) ^ (offset >> 6)) & 1; + offset ^= (difference << 5) | (difference << 6); + } + return offset; } uvec4 ReadTexel(uint offset) { diff --git a/src/video_core/host_shaders/block_linear_unswizzle_3d.comp b/src/video_core/host_shaders/block_linear_unswizzle_3d.comp index 34e4f93307..8dfaf0cd1a 100644 --- a/src/video_core/host_shaders/block_linear_unswizzle_3d.comp +++ b/src/video_core/host_shaders/block_linear_unswizzle_3d.comp @@ -41,6 +41,7 @@ UNIFORM(6) uint block_height; UNIFORM(7) uint block_height_mask; UNIFORM(8) uint block_depth; UNIFORM(9) uint block_depth_mask; +UNIFORM(10) uint native_tiling; END_PUSH_CONSTANTS #if HAS_EXTENDED_TYPES @@ -79,7 +80,12 @@ uint SwizzleTable(uint pos) { uint SwizzleOffset(uvec2 pos) { pos = pos & SWIZZLE_MASK; - return SwizzleTable(pos.y * 64 + pos.x); + uint offset = SwizzleTable(pos.y * 64 + pos.x); + if (native_tiling != 0) { + uint difference = ((offset >> 5) ^ (offset >> 6)) & 1; + offset ^= (difference << 5) | (difference << 6); + } + return offset; } uvec4 ReadTexel(uint offset) { diff --git a/src/video_core/host_shaders/block_linear_unswizzle_3d_bcn.comp b/src/video_core/host_shaders/block_linear_unswizzle_3d_bcn.comp index e084837f25..28b1dcf599 100644 --- a/src/video_core/host_shaders/block_linear_unswizzle_3d_bcn.comp +++ b/src/video_core/host_shaders/block_linear_unswizzle_3d_bcn.comp @@ -42,7 +42,7 @@ layout(push_constant) uniform PushConstants { uint block_depth; // Offset 48 uint block_depth_mask; // Offset 52 - int _pad; // Offset 56 + uint native_tiling; // Offset 56 ivec3 destination; // Offset 60 } pc; @@ -59,6 +59,7 @@ BEGIN_PUSH_CONSTANTS UNIFORM(8) uint block_depth; UNIFORM(9) uint block_depth_mask; UNIFORM(10) uvec3 blocks_dim; + UNIFORM(11) uint native_tiling; END_PUSH_CONSTANTS #define pc // Map pc prefix to nothing for OpenGL compatibility #endif @@ -103,7 +104,12 @@ uint SwizzleTable(uint pos) { // --- Helpers --- uint SwizzleOffset(uvec2 pos) { pos &= SWIZZLE_MASK; - return SwizzleTable(pos.y * 64u + pos.x); + uint offset = SwizzleTable(pos.y * 64 + pos.x); + if (pc.native_tiling != 0) { + uint difference = ((offset >> 5) ^ (offset >> 6)) & 1; + offset ^= (difference << 5) | (difference << 6); + } + return offset; } uvec4 ReadTexel(uint offset) { diff --git a/src/video_core/macro.cpp b/src/video_core/macro.cpp index 86f528af15..a8acd81682 100644 --- a/src/video_core/macro.cpp +++ b/src/video_core/macro.cpp @@ -4,6 +4,8 @@ // SPDX-FileCopyrightText: Copyright 2020 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include +#include #include #include #include @@ -1339,6 +1341,34 @@ static void Dump(u64 hash, std::span code, bool decompiled = false) { void MacroEngine::Execute(Core::System& system, Engines::Maxwell3D& maxwell3d, u32 method, std::span parameters) { + if (native_turing) { + maxwell3d.RefreshParameters(); + const auto context = reinterpret_cast(this); + static std::unordered_set captured_contexts; + if (std::getenv("EDEN_NATIVE_FRAME_CAPTURE") && captured_contexts.insert(context).second) { + std::ofstream dump("native-graphics-mme-" + std::to_string(context) + ".bin", std::ios::binary); + dump.write(reinterpret_cast(native_code.data()), sizeof(native_code)); + } + LOG_DEBUG(HW_GPU, "Native MME context={:#x}", context); + LOG_DEBUG(HW_GPU, "Native MME start={:#x}, parameters={}, data={:08x}", method, parameters.size(), fmt::join(parameters, ",")); + const auto result = native_interpreter.Execute(native_code, method, parameters, + [&](u32 address) { return address < Maxwell3D::Regs::NUM_REGS ? maxwell3d.GetRegisterValue(address) : 0U; }, + [&](u32 address, u32 value) { + if (address >= 0xe00) return false; + maxwell3d.CallMethod(system, address, value, true); + return true; + }); + if (!result.success) { + native_execution_failed = true; + if (std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS")) { + std::ofstream dump("native-mme-" + std::to_string(method) + ".bin", std::ios::binary); + dump.write(reinterpret_cast(native_code.data()), sizeof(native_code)); + } + LOG_ERROR(HW_GPU, "Native MME failed: start={:#x}, pc={:#x}, consumed={}/{}, error={}", + method, result.pc, result.consumed, parameters.size(), result.error); + } + return; + } const auto execute_variant = [&system, &maxwell3d, ¶meters, method](AnyCachedMacro& cached) { if (std::holds_alternative(cached) || diff --git a/src/video_core/macro.h b/src/video_core/macro.h index 90519a5429..52dfe7c82c 100644 --- a/src/video_core/macro.h +++ b/src/video_core/macro.h @@ -7,6 +7,7 @@ #pragma once #include +#include "video_core/macro_turing.h" #include #include #include @@ -223,10 +224,15 @@ struct MacroEngine { MacroEngine(bool is_interpreted_) noexcept : is_interpreted{is_interpreted_} {} // Store the uploaded macro code to compile them when they're called. inline void AddCode(u32 method, u32 data) noexcept { + if (native_turing) { + if (native_write_offset < native_code.size()) native_code[native_write_offset++] = data; + return; + } uploaded_macro_code[method].push_back(data); } // Clear the code associated with a method. inline void ClearCode(u32 method) noexcept { + if (native_turing) { native_write_offset = method < 0x1000 ? method * 3 : static_cast(native_code.size()); return; } macro_cache.erase(method); uploaded_macro_code.erase(method); } @@ -240,6 +246,19 @@ struct MacroEngine { ::Common::unordered_map macro_cache; ::Common::unordered_map> uploaded_macro_code; bool is_interpreted; + bool native_turing{}; + bool native_execution_failed{}; + u32 native_write_offset{}; + std::array native_code{}; + TuringMacroInterpreter native_interpreter; + +}; + +struct NativeMacroContext { + MacroEngine engine{true}; + std::array positions{}; + MacroEngine compute_engine{true}; + std::array compute_positions{}; }; } // namespace Tegra diff --git a/src/video_core/macro_turing.h b/src/video_core/macro_turing.h new file mode 100644 index 0000000000..fe82ce4a4f --- /dev/null +++ b/src/video_core/macro_turing.h @@ -0,0 +1,213 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later +// Instruction semantics adapted from Mesa mme_tu104_sim.c and mme_tu104.c: +// Copyright © 2022 Collabora Ltd. SPDX-License-Identifier: MIT +#pragma once + +#include +#include +#include +#include +#include +#include + +namespace Tegra { + +// TU104/GA10x MME has two ALUs and two outputs per 96-bit instruction. +class TuringMacroInterpreter { +public: + using Word = std::uint32_t; + struct Result { + bool success; + Word pc; + std::size_t consumed; + std::string_view error; + }; + + template + Result Execute(std::span code, Word start, + std::span parameters, Read read, Emit emit) { + std::array regs{}; + Word initialized{}, carry{}, method{}, increment{}; + Word pc = start, loop_count{}, loop_start{}, loop_end{}; + std::size_t consumed{}; + bool end_next{}, stopped{}; + std::string_view error; + const auto fail = [&](std::string_view message) { if (error.empty()) error = message; }; + for (Word steps = 0; steps < 1000000; ++steps) { + if (static_cast(pc) * 3 + 2 >= code.size()) { + fail("instruction outside uploaded RAM"); + break; + } + const std::array bits{code[pc * 3 + 2], code[pc * 3 + 1], code[pc * 3]}; + const auto field = [&](Word low, Word count) { + std::uint64_t value = bits[low / 32]; + if (low / 32 < 2) value |= std::uint64_t{bits[low / 32 + 1]} << 32; + return Word((value >> (low % 32)) & ((std::uint64_t{1} << count) - 1)); + }; + const std::array imm{field(30, 16), field(66, 16)}; + const std::array op{field(10, 5), field(46, 5)}; + const std::array dst{field(15, 5), field(51, 5)}; + const std::array src_a{field(20, 5), field(56, 5)}; + const std::array src_b{field(25, 5), field(61, 5)}; + const std::array out_method{field(82, 3), field(89, 3)}; + const std::array out_emit{field(85, 4), field(92, 4)}; + const Word pred_reg = field(5, 5), pred_mode = field(1, 4); + std::array load{}, alu{}; + for (Word i = 0; i < 2; ++i) { + const Word reg = 28 + i, out = 3 + i; + const bool needed = pred_reg == reg || src_a[0] == reg || src_a[1] == reg || + src_b[0] == reg || src_b[1] == reg || out_method[0] == out || + out_method[1] == out || out_emit[0] == out || out_emit[1] == out; + if (needed) { + if (consumed >= parameters.size()) { fail("parameter FIFO underflow"); break; } + load[i] = parameters[consumed++]; + } + } + // Both ALUs read the register file before either destination is + // committed. Carry chaining remains explicit between the lanes. + const auto input_regs = regs; + const Word input_initialized = initialized; + const auto get = [&](Word reg, Word lane) -> Word { + if (reg < 24) { + if (!(input_initialized & (1U << reg))) fail("uninitialized register read"); + return input_regs[reg]; + } + switch (reg) { + case 24: return 0; + case 25: return Word(std::int32_t(std::int16_t(imm[lane]))); + case 26: return Word(std::int32_t(std::int16_t(imm[1 - lane]))); + case 27: return (imm[0] << 16) | imm[1]; + case 28: return load[0]; + case 29: return load[1]; + default: fail("invalid source register"); return 0; + } + }; + constexpr std::array predicates{ + "UUUU", "TTTT", "FFFF", "TTUU", "FFUU", "TFUU", "TUUU", "FUUU", + "UUTT", "UUTF", "UUTU", "UUFT", "UUFF", "UUFU", "UUUT", "UUUF"}; + Word mask = 15; + if (pred_mode) { + const char skipped = get(pred_reg, 0) ? 'T' : 'F'; + mask = 0; + for (Word i = 0; i < 4; ++i) + if (predicates[pred_mode][i] != skipped) mask |= 1U << i; + } + Word next_pc = pc + 1; + bool branch_cancels_end{}; + for (Word lane = 0; lane < 2 && error.empty(); ++lane) { + if (!(mask & (1U << lane))) continue; + const Word x = get(src_a[lane], lane), y = get(src_b[lane], lane); + Word value{}; + const auto condition = [&](Word comparison) { + switch (comparison) { + case 17: case 25: return std::int32_t(x) < std::int32_t(y); + case 18: case 26: return x < y; + case 19: case 27: return std::int32_t(x) <= std::int32_t(y); + case 20: case 28: return x <= y; + default: return x == y; + } + }; + switch (op[lane]) { + case 0: value = x + y; carry = value < x; break; + case 1: value = x + y + carry; break; + case 2: value = x - y; carry = value > x; break; + case 3: value = x - y - carry; break; + case 4: case 6: { + const std::uint64_t a = op[lane] == 4 ? std::uint64_t(std::int64_t(std::int32_t(x))) : x; + const std::uint64_t b = op[lane] == 4 ? std::uint64_t(std::int64_t(std::int32_t(y))) : y; + const auto product = a * b; + value = Word(product); carry = Word(product >> 32); break; + } + case 5: value = carry; break; + case 7: + // EXTENDED: x = 0x1000 is the load barrier (Mesa mme_tu104_load_barrier), + // nothing to do in a sequential interpreter. + if (x != 0x1000) fail("unsupported extended operation"); + break; + case 8: value = Word(std::countl_zero(x)); break; + case 9: value = x << (y & 31); break; + case 10: value = x >> (y & 31); break; + case 11: value = Word(std::int32_t(x) >> (y & 31)); break; + case 12: value = x & y; break; + case 13: value = ~(x & y); break; + case 14: value = x | y; break; + case 15: value = x ^ y; break; + case 16: { + const Word dest = (imm[lane] >> 10) & 63, count = (imm[lane] >> 5) & 31; + const Word source = imm[lane] & 31; + if (dest >= 32) { fail("invalid merge destination"); break; } + const Word bits_mask = (1U << count) - 1; + value = (x & ~(bits_mask << dest)) | (((y >> source) & bits_mask) << dest); + break; + } + case 17: case 18: case 19: case 20: case 21: + value = condition(op[lane]) ? ~0U : 0; break; + case 22: value = read((x + y) & 0xffff); break; + case 23: + if (loop_count || imm[lane] < 2) { fail("invalid loop"); break; } + // The native SDK emits zero-count loops around FIFO reads. + // Unlike Mesa's test simulator, these must skip the body. + if (x == 0) { next_pc = pc + imm[lane]; break; } + loop_count = x - 1; loop_start = pc; loop_end = pc + imm[lane] - 1; + break; + case 24: next_pc = pc + (imm[lane] & 0x7fff); break; + case 25: case 26: case 27: case 28: case 29: + if (condition(op[lane]) == bool(imm[lane] & 0x8000)) { + const auto offset = std::int32_t(imm[lane] << 19) >> 19; + if (offset == -4096) stopped = true; + else { + next_pc = Word(std::int32_t(pc) + offset); + // Native back edges use bit 14 to suppress the + // fallthrough termination when the branch is taken. + branch_cancels_end = (imm[lane] & 0x4000) != 0; + } + } + break; + case 30: + if (x >= dram.size()) { fail("data RAM read out of bounds"); break; } + value = dram[x]; break; + case 31: + if (x >= dram.size()) { fail("data RAM write out of bounds"); break; } + dram[x] = y; break; + } + alu[lane] = value; + if (dst[lane] < 24) { regs[dst[lane]] = value; initialized |= 1U << dst[lane]; } + else if (dst[lane] != 24) fail("invalid destination register"); + } + const auto output = [&](Word selector) -> Word { + switch (selector) { + case 1: return alu[0]; case 2: return alu[1]; + case 3: return load[0]; case 4: return load[1]; + case 5: return imm[0]; case 6: return imm[1]; + case 8: return imm[0] >> 12; case 9: return imm[1] >> 12; + case 10: return (imm[0] << 16) | imm[1]; + default: fail("invalid output selector"); return 0; + } + }; + for (Word lane = 0; lane < 2 && error.empty(); ++lane) { + if (!(mask & (4U << lane))) continue; + if (out_method[lane]) { + const Word address = output(out_method[lane]); + method = address & 0xfff; increment = (address >> 12) & 15; + } + if (out_emit[lane]) { + const Word value = output(out_emit[lane]); + if (error.empty() && !emit(method, value)) fail("unsupported emitted method"); + method += increment; + } + } + if (!error.empty()) break; + if (end_next || stopped) return {true, pc, consumed, {}}; + end_next = field(0, 1) != 0 && !branch_cancels_end; + if (loop_count && pc == loop_end) { --loop_count; next_pc = loop_start + 1; } + pc = next_pc; + } + if (error.empty()) error = "instruction limit exceeded"; + return {false, pc, consumed, error}; + } + + std::array dram{}; +}; + +} // namespace Tegra diff --git a/src/video_core/memory_manager.cpp b/src/video_core/memory_manager.cpp index 52e9db189d..9a9fdab7bc 100644 --- a/src/video_core/memory_manager.cpp +++ b/src/video_core/memory_manager.cpp @@ -22,6 +22,10 @@ #include "video_core/texture_cache/util.h" namespace Tegra { +bool MemoryManager::IsShuttingDown() const { + return system.IsShuttingDown(); +} + using Tegra::Memory::GuestMemoryFlags; std::atomic MemoryManager::unique_identifier_generator{}; diff --git a/src/video_core/memory_manager.h b/src/video_core/memory_manager.h index 1aaf2c1086..105cb14e23 100644 --- a/src/video_core/memory_manager.h +++ b/src/video_core/memory_manager.h @@ -47,6 +47,15 @@ public: void SetNativeSyncpointShim(GPUVAddr address); static constexpr u32 NumShimSyncpoints = 192; // Tegra::Host1x::SyncpointManager + /// Switch 2 lays out block-linear surfaces with the console's GOB swizzle. + void EnableNativeTiling() { + native_tiling = true; + } + bool UsesNativeTiling() const { + return native_tiling; + } + bool IsShuttingDown() const; + static constexpr bool HAS_FLUSH_INVALIDATION = true; inline size_t GetID() const noexcept { @@ -225,6 +234,7 @@ private: boost::container::small_vector, 32> page_stash{}; boost::container::small_vector, 32> page_stash2{}; + bool native_tiling{}; mutable std::mutex guard; static constexpr size_t continuous_bits = 64; diff --git a/src/video_core/native_qmd.h b/src/video_core/native_qmd.h new file mode 100644 index 0000000000..c9ae3b8cc0 --- /dev/null +++ b/src/video_core/native_qmd.h @@ -0,0 +1,112 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later +#pragma once +#include +#include +#include +#include +#include + +namespace Tegra { + +// Native QMD 3.0 dependency scheduling; nonempty grids require a dispatch backend. +class NativeQmdScheduler { +public: + using Word = std::uint32_t; + using Address = std::uint64_t; + using Descriptor = std::array; + struct RejectShaderDispatch { bool operator()(const Descriptor&) const { return false; } }; + template + std::string_view Signal(Address address, Word argument, Read read, Release release) { + return Signal(address, argument, read, release, RejectShaderDispatch{}); + } + template + std::string_view Signal(Address address, Word argument, Read read, Release release, Dispatch dispatch) { + struct Pending { Address address; Word action; Word count; }; + std::array pending{}; + std::size_t count = 1; + pending[0] = {address, argument & 15, ((argument >> 14) & 0x3ff) + 1}; + for (std::size_t i = 0; i < count; ++i) { + const auto request = pending[i]; + if (states.size() >= 4096 && !states.contains(request.address)) return "QMD cache limit"; + auto& state = states[request.address]; + if (request.action == 3 || request.action == 2) { + if (request.action == 3 || !state.loaded) { + if (!read(request.address, state.descriptor)) return "unmapped QMD"; + const auto& qmd = state.descriptor; + if ((qmd[18] & 0xff) != 0x30 || (qmd[4] & 0x100)) return "unsupported QMD version or queue"; + if constexpr (std::is_same_v) { + if (qmd[12] && (qmd[13] & 0xffff) && (qmd[14] & 0xffff)) + return "native shader dispatch required"; + } + state.dependencies = std::int32_t((qmd[4] >> 14) & 0xffff) + state.deferred; + state.deferred = 0; + state.loaded = true; + state.completed = false; + } + state.scheduled = true; + } else if (request.action == 7 || request.action == 12) { + const auto delta = request.action == 7 ? -std::int32_t(request.count) : std::int32_t(request.count); + if (!state.loaded || state.completed) { state.deferred += delta; continue; } + state.dependencies += delta; + } else if (request.action == 0 || request.action == 1) { + continue; + } else { + return "unsupported PCAS action"; + } + if (!state.loaded || !state.scheduled || state.completed || state.dependencies > 0) continue; + if (state.dependencies < 0) return "QMD dependency underflow"; + // The SDK patches release payloads while a scheduled QMD is + // waiting for its remaining dependencies. Fetch them at execution. + if (!read(request.address, state.descriptor)) return "unmapped QMD"; + const auto qmd = state.descriptor; + if ((qmd[18] & 0xff) != 0x30 || (qmd[4] & 0x100)) return "unsupported QMD version or queue"; + // Validate both release slots before making either visible. + for (Word slot = 0; slot < 2; ++slot) { + const Word flags = qmd[25 + slot * 4]; + if (!(flags & (1U << 23))) continue; + const bool add = (flags & (1U << 26)) != 0; + if ((flags & (0x3bU << 24)) || (flags >> 30) != 1 || + (add && (flags & (7U << 20)))) return "unsupported QMD release format"; + if (add) { + if constexpr (!std::is_invocable_r_v) + return "unsupported QMD reduction callback"; + } + } + if (qmd[12] && (qmd[13] & 0xffff) && (qmd[14] & 0xffff) && !dispatch(qmd)) + return "native shader dispatch required"; + state.completed = true; + for (Word slot = 0; slot < 2; ++slot) { + const Word flags = qmd[25 + slot * 4]; + if (!(flags & (1U << 23))) continue; + const Address destination = (Address{flags & 0xff} << 32) | qmd[24 + slot * 4]; + bool released{}; + if constexpr (std::is_invocable_r_v) { + released = release(destination, qmd[26 + slot * 4], flags); + } else { + released = release(destination, qmd[26 + slot * 4]); + } + if (!released) return "unmapped QMD release"; + } + if (qmd[16] & 1) { + const Word action = (qmd[16] >> 1) & 7; + if (action != 4 && action != 3 && action != 1) return "unsupported dependent QMD action"; + if (count == pending.size()) return "QMD chain limit"; + pending[count++] = {Address{qmd[15]} << 8, action == 4 ? 7U : action == 3 ? 3U : 2U, 1}; + } + if (qmd[16] & (1U << 5)) return "second dependent QMD unsupported"; + } + return {}; + } +private: + struct State { + Descriptor descriptor{}; + std::int32_t dependencies{}; + std::int32_t deferred{}; + bool loaded{}; + bool scheduled{}; + bool completed{}; + }; + std::unordered_map states; +}; +} // namespace Tegra diff --git a/src/video_core/native_query_trace.h b/src/video_core/native_query_trace.h new file mode 100644 index 0000000000..eae258dc04 --- /dev/null +++ b/src/video_core/native_query_trace.h @@ -0,0 +1,42 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later +#pragma once + +#include +#include +#include +#include +#include "common/common_types.h" + +namespace VideoCommon { +namespace NativeQueryDiagnostics { +struct State { + State() : file{"native-query-trace.csv", std::ios::trunc} { + file << "event,asid,gpu,cpu,type,value,flags\n"; + } + std::mutex mutex; + std::ofstream file; + size_t records{}; +}; +inline State& GetState() { + static State state; + return state; +} +} // namespace NativeQueryDiagnostics + +// Events: 0=request, 1=written, 2=invalidated, 3=unsynchronized, +// 4=channel suspended, 5=channel resumed, 6=buffer readback, +// 7/8=QMD release slots observed, 9=PCAS request, +// 10=DMA destination (value=source, flags=line bytes), 11=inline write. +// Keep this opt-in and bounded. +inline void TraceNativeQuery(u32 event, u64 asid, u64 gpu, u64 cpu, u32 type, u64 value, u32 flags) { + if (!std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS")) return; + auto& state = NativeQueryDiagnostics::GetState(); + const std::scoped_lock lock{state.mutex}; + if (state.records++ < 1000000 || event == 4 || event == 5) { + state.file << event << ',' << asid << ',' << std::hex << gpu << ',' << cpu << ',' + << type << ',' << value << ',' << flags << std::dec << '\n'; + } + if (event == 4 || event == 5) state.file.flush(); +} +} // namespace VideoCommon diff --git a/src/video_core/renderer_vulkan/vk_texture_cache.cpp b/src/video_core/renderer_vulkan/vk_texture_cache.cpp index 583cb6edaf..29afaac01b 100644 --- a/src/video_core/renderer_vulkan/vk_texture_cache.cpp +++ b/src/video_core/renderer_vulkan/vk_texture_cache.cpp @@ -2085,6 +2085,11 @@ void Image::DownloadMemory(VkBuffer buffer, size_t offset, void Image::DownloadMemory(std::span buffers_span, std::span offsets_span, std::span copies) { + if (!original_image) { + LOG_CRITICAL(Render_Vulkan, "Missing image at download: gpu={:#x}, cpu={:#x}, type={}, format={}, flags={:#x}, size={}, dimensions={}x{}", + gpu_addr, cpu_addr, static_cast(info.type), static_cast(info.format), + static_cast(flags), unswizzled_size_bytes, info.size.width, info.size.height); + } const bool is_rescaled = True(flags & ImageFlagBits::Rescaled); if (is_rescaled) { ScaleDown(); diff --git a/src/video_core/shader_cache.cpp b/src/video_core/shader_cache.cpp index f1aff2810f..75e5432ab2 100644 --- a/src/video_core/shader_cache.cpp +++ b/src/video_core/shader_cache.cpp @@ -11,6 +11,7 @@ #include "common/assert.h" #include "shader_recompiler/frontend/maxwell/control_flow.h" +#include "shader_recompiler/exception.h" #include "shader_recompiler/object_pool.h" #include "video_core/control/channel_state.h" #include "video_core/dirty_flags.h" @@ -43,7 +44,7 @@ void ShaderCache::SyncGuestHost() { ShaderCache::ShaderCache(Tegra::MaxwellDeviceMemoryManager& device_memory_) : device_memory{device_memory_} {} -bool ShaderCache::RefreshStages(std::array& unique_hashes) { +bool ShaderCache::RefreshStages(std::array& unique_hashes) try { auto& dirty{maxwell3d->dirty.flags}; if (!dirty[VideoCommon::Dirty::Shaders]) { return last_shaders_valid; @@ -51,6 +52,7 @@ bool ShaderCache::RefreshStages(std::array& unique_hashes) { dirty[VideoCommon::Dirty::Shaders] = false; const GPUVAddr base_addr{maxwell3d->regs.program_region.Address()}; + const bool native = maxwell3d->GetNativeMacroEngine().native_turing; for (size_t index = 0; index < Tegra::Engines::Maxwell3D::Regs::MaxShaderProgram; ++index) { if (!maxwell3d->regs.IsShaderConfigEnabled(index)) { unique_hashes[index] = 0; @@ -63,7 +65,9 @@ bool ShaderCache::RefreshStages(std::array& unique_hashes) { unique_hashes[index] = 0; continue; } - const GPUVAddr shader_addr{base_addr + shader_config.offset}; + const GPUVAddr shader_addr{native + ? (static_cast(shader_config.reserved[0] & 0xff) << 32) | shader_config.reserved[1] + : base_addr + shader_config.offset}; const std::optional cpu_shader_addr{gpu_memory->GpuToCpuAddress(shader_addr)}; if (!cpu_shader_addr) { LOG_ERROR(HW_GPU, "Invalid GPU address for shader {:#016x}", shader_addr); @@ -72,8 +76,8 @@ bool ShaderCache::RefreshStages(std::array& unique_hashes) { } const ShaderInfo* shader_info{TryGet(*cpu_shader_addr)}; if (!shader_info) { - const u32 start_address{shader_config.offset}; - GraphicsEnvironment env{*maxwell3d, *gpu_memory, program, base_addr, start_address}; + const u32 start_address{native ? 0 : shader_config.offset}; + GraphicsEnvironment env{*maxwell3d, *gpu_memory, program, native ? shader_addr : base_addr, start_address}; shader_info = MakeShaderInfo(env, *cpu_shader_addr); } shader_infos[index] = shader_info; @@ -81,6 +85,13 @@ bool ShaderCache::RefreshStages(std::array& unique_hashes) { } last_shaders_valid = true; return true; +} catch (const Shader::Exception& error) { + auto& native = maxwell3d->GetNativeMacroEngine(); + if (!native.native_turing) throw; + native.native_execution_failed = true; + last_shaders_valid = false; + LOG_ERROR(HW_GPU, "Native graphics shader rejected: {}", error.what()); + return false; } const ShaderInfo* ShaderCache::ComputeShader() { @@ -103,14 +114,19 @@ void ShaderCache::GetGraphicsEnvironments(GraphicsEnvironments& result, const std::array& unique_hashes) { size_t env_index{}; const GPUVAddr base_addr{maxwell3d->regs.program_region.Address()}; + const bool native = maxwell3d->GetNativeMacroEngine().native_turing; for (size_t index = 0; index < NUM_PROGRAMS; ++index) { if (unique_hashes[index] == 0) { continue; } const auto program{static_cast(index)}; auto& env{result.envs[index]}; - const u32 start_address{maxwell3d->regs.pipelines[index].offset}; - env = GraphicsEnvironment{*maxwell3d, *gpu_memory, program, base_addr, start_address}; + const auto& config = maxwell3d->regs.pipelines[index]; + const u32 start_address{native ? 0 : config.offset}; + const GPUVAddr address = native + ? (static_cast(config.reserved[0] & 0xff) << 32) | config.reserved[1] + : base_addr; + env = GraphicsEnvironment{*maxwell3d, *gpu_memory, program, address, start_address}; env.SetCachedSize(shader_infos[index]->size_bytes); result.env_ptrs[env_index++] = &env; } @@ -241,7 +257,7 @@ const ShaderInfo* ShaderCache::MakeShaderInfo(GenericEnvironment& env, VAddr cpu info->unique_hash = env.CalculateHash(); info->size_bytes = env.ReadSizeBytes(); } - const size_t size_bytes{info->size_bytes}; + const size_t size_bytes{env.SourceSizeBytes() ? env.SourceSizeBytes() : info->size_bytes}; const ShaderInfo* const result{info.get()}; Register(std::move(info), cpu_addr, size_bytes); return result; diff --git a/src/video_core/shader_environment.cpp b/src/video_core/shader_environment.cpp index 4d5db7e511..fc2a26a568 100644 --- a/src/video_core/shader_environment.cpp +++ b/src/video_core/shader_environment.cpp @@ -5,7 +5,11 @@ // SPDX-License-Identifier: GPL-2.0-or-later #include +#include +#include +#include #include +#include #include #include #include @@ -21,6 +25,11 @@ #include "common/settings.h" #include #include "shader_recompiler/environment.h" +#include "shader_recompiler/exception.h" +#include "shader_recompiler/frontend/ampere/lower_to_maxwell.h" +#include "shader_recompiler/frontend/ampere/program_override.h" +#include "shader_recompiler/frontend/ampere/program_capture.h" +#include "shader_recompiler/frontend/ampere/program_address_map.h" #include "video_core/engines/kepler_compute.h" #include "video_core/memory_manager.h" #include "video_core/shader_environment.h" @@ -149,11 +158,26 @@ u64 GenericEnvironment::ReadInstruction(u32 address) { if (address >= cached_lowest && address < cached_highest) { return code[(address - cached_lowest) / INST_SIZE]; } + if (native_source_size) { + throw Shader::RuntimeError("Lowered native instruction address out of range: {:#x}", address); + } has_unbound_instructions = true; return gpu_memory->Read(program_base + address); } +u32 GenericEnvironment::TranslateProgramAddress(u32 address) const { + if (const auto translated = Shader::Ampere::NativeProgramAddress(code, address)) return *translated; + if (native_source_size || Shader::Ampere::HasNativeAddressMap(code)) + throw Shader::RuntimeError("Invalid native indirect branch address {:#x}", address); + return address; +} + std::optional GenericEnvironment::Analyze() { + if (native_source_size) { + // Include the converted header and a version tag to isolate this experimental ABI. + return Common::CityHash64(reinterpret_cast(code.data()), + code.size() * sizeof(u64)) ^ 0x4f554e4345000002ULL; + } const std::optional size{TryFindSize()}; if (!size) { return std::nullopt; @@ -164,6 +188,12 @@ std::optional GenericEnvironment::Analyze() { } void GenericEnvironment::SetCachedSize(size_t size_bytes) { + if (native_source_size) { + if (size_bytes != CachedSizeBytes()) { + throw Shader::RuntimeError("Native shader changed while constructing pipeline"); + } + return; + } cached_lowest = start_address; cached_highest = start_address + static_cast(size_bytes); code.resize(CachedSizeWords()); @@ -291,6 +321,125 @@ Tegra::Texture::TICEntry GenericEnvironment::ReadTextureInfo(GPUVAddr tic_addr, return entry; } +void GenericEnvironment::InitializeNativeCode(u32 source_header_size) { + using Shader::Ampere::Instruction; + const u32 native_header_size = source_header_size; + constexpr u32 maximum_instructions = Shader::Ampere::ProgramCaptureState::MaximumInstructions; + std::vector native; + Shader::Ampere::ProgramCaptureState capture; + for (u32 index = 0; index < maximum_instructions; ++index) { + const GPUVAddr address = program_base + native_header_size + index * sizeof(Instruction); + if (!gpu_memory->IsFullyMappedRange(address, sizeof(Instruction))) { + // The program may end before a branch target past its last EXIT (the capture can't + // tell); stop at the mapping's end and let the lowering judge what was captured. + if (!native.empty()) break; + throw Shader::RuntimeError("Unmapped native shader instruction at {:#x}", address); + } + auto& instruction = native.emplace_back(); + gpu_memory->ReadBlock(address, &instruction, sizeof(instruction)); + const auto status = capture.Observe(instruction, index); + if (status == Shader::Ampere::ProgramCaptureState::Result::InvalidTarget) + throw Shader::RuntimeError("Native shader branch outside capture bounds"); + if (status == Shader::Ampere::ProgramCaptureState::Result::Complete) break; + } + auto lowered = Shader::Ampere::LowerToMaxwell(native); + // EDEN_NX2_SHADER_DUMP=

: keep every captured native program (named by its key) so the + // lowering can be tested offline against all of a game's shaders. + if (static const char* const dump_directory = std::getenv("EDEN_NX2_SHADER_DUMP"); + dump_directory != nullptr) { + const auto folder = std::filesystem::path{dump_directory}; + const auto path = folder / fmt::format("{:016x}-{}-{:x}.bin", Shader::Ampere::NativeProgramKey(native), + source_header_size == 0 ? "cs" : "gfx", program_base); + if (!std::filesystem::exists(path)) { + std::filesystem::create_directories(folder); + std::ofstream dump(path, std::ios::binary); + dump.write(reinterpret_cast(native.data()), + native.size() * sizeof(Shader::Ampere::Instruction)); + } + } + if (!lowered) { + if (const char* directory = std::getenv("EDEN_NATIVE_SHADER_OVERRIDE_DIR")) { + const auto key = fmt::format("{:016x}", Shader::Ampere::NativeProgramKey(native)); + const auto folder = std::filesystem::path{directory}; + std::filesystem::create_directories(folder); + const auto source_path = folder / (key + ".source.bin"); + const auto override_path = folder / (key + ".override.bin"); + { + std::ofstream dump(source_path, std::ios::binary); + dump.write(reinterpret_cast(native.data()), native.size() * sizeof(Shader::Ampere::Instruction)); + } + { + // Retain mapped bytes beyond the first EXIT for investigating + // native CALL targets in compiler library subroutines. + std::ofstream mapped(folder / (key + ".mapped.bin"), std::ios::binary); + std::array block{}; + for (u32 offset = 0; offset < 0x10000; offset += 0x1000) { + if (!gpu_memory->IsFullyMappedRange(program_base + offset, block.size())) break; + gpu_memory->ReadBlock(program_base + offset, block.data(), block.size()); + mapped.write(reinterpret_cast(block.data()), block.size()); + } + } + LOG_WARNING(Shader, "Native shader paused: key={}, offset={:#x}, reason={}, source={}", + key, lowered.error_offset, lowered.error, source_path.string()); + const auto deadline = std::chrono::steady_clock::now() + std::chrono::minutes{30}; + while (!gpu_memory->IsShuttingDown() && std::chrono::steady_clock::now() < deadline) { + if (std::filesystem::exists(override_path)) { + std::ifstream file(override_path, std::ios::binary | std::ios::ate); + const auto length = file.tellg(); + if (length <= 0 || length > 16 * 1024 * 1024) { + throw Shader::RuntimeError("Invalid native shader override size"); + } + std::vector bytes(static_cast(length)); + file.seekg(0); + file.read(reinterpret_cast(bytes.data()), bytes.size()); + if (!file) throw Shader::RuntimeError("Native shader override read failed"); + const u32 source_bytes = Shader::Ampere::ProgramOverrideSourceSize(bytes); + if (source_bytes > native.size() * sizeof(Shader::Ampere::Instruction)) { + if (!gpu_memory->IsFullyMappedRange(program_base + native_header_size, source_bytes)) + throw Shader::RuntimeError("Unmapped extended native shader source"); + native.resize(source_bytes / sizeof(Shader::Ampere::Instruction)); + gpu_memory->ReadBlock(program_base + native_header_size, native.data(), source_bytes); + } + lowered = Shader::Ampere::DecodeProgramOverride(bytes, native); + if (!lowered) throw Shader::RuntimeError("Native shader override rejected: {}", lowered.error); + LOG_WARNING(Shader, "Native shader resumed: key={}, instructions={}", key, lowered.instructions.size()); + break; + } + std::this_thread::sleep_for(std::chrono::milliseconds{100}); + } + } + } + if (!lowered) { + const size_t failed = lowered.error_offset / sizeof(Shader::Ampere::Instruction); + const auto words = failed < native.size() ? native[failed].words : std::array{}; + throw Shader::RuntimeError( + "Native shader at {:#x}, instruction {:#x}: {} [op {:03x} {:08x} {:08x} {:08x} {:08x}]", + program_base, lowered.error_offset, lowered.error, words[0] & 0xfff, words[0], + words[1], words[2], words[3]); + } + if (lowered.required_texture_buffer != ~0U) { + // Native TEX carries its cbuf index explicitly; TEXS obtains it from the environment. + texture_bound = lowered.required_texture_buffer; + } + native_source_size = native_header_size + lowered.source_size; + code.resize(sizeof(sph) / sizeof(u64)); + std::memcpy(code.data(), &sph, sizeof(sph)); + for (const u64 instruction : lowered.instructions) { + if ((code.size() * sizeof(u64)) % 32 == 0) { + code.push_back(0); // Maxwell scheduling word, absent from native instructions. + } + code.push_back(instruction); + } + // Keep the original padding ABI, then retain the native address map as + // unreachable data. Existing disk-cache serialization carries it with code. + code.push_back(0); + Shader::Ampere::AppendNativeAddressMap(code, lowered.source_instruction_offsets); + cached_lowest = 0; + cached_highest = static_cast((code.size() - 1) * sizeof(u64)); + LOG_INFO(Shader, "Lowered native shader at {:#x}: {} source bytes, {} instructions", + program_base, native_source_size, lowered.instructions.size()); +} + GraphicsEnvironment::GraphicsEnvironment(Tegra::Engines::Maxwell3D& maxwell3d_, Tegra::MemoryManager& gpu_memory_, Maxwell::ShaderType program, GPUVAddr program_base_, @@ -298,6 +447,23 @@ GraphicsEnvironment::GraphicsEnvironment(Tegra::Engines::Maxwell3D& maxwell3d_, : GenericEnvironment{gpu_memory_, program_base_, start_address_}, maxwell3d{&maxwell3d_} { gpu_memory->ReadBlock(program_base + start_address, &sph, sizeof(sph)); initial_offset = sizeof(sph); + texture_bound = maxwell3d->regs.bindless_texture_const_buffer_slot; + if (maxwell3d->GetNativeMacroEngine().native_turing) { + if (program != Maxwell::ShaderType::VertexB && program != Maxwell::ShaderType::Pixel) { + throw Shader::RuntimeError("Unsupported native graphics shader stage {}", static_cast(program)); + } + if (std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS")) { + const u32 native_stage = program == Maxwell::ShaderType::Pixel ? 4 : 0; + const auto& table = maxwell3d->state.shader_stages[native_stage].const_buffers[1]; + if (table.enabled && table.size && table.size <= Tegra::Engines::Maxwell3D::Regs::MaxConstBufferSize) { + std::vector bytes(table.size); + gpu_memory->ReadBlock(table.address, bytes.data(), bytes.size()); + std::ofstream dump(fmt::format("native-graphics-cbuf-{}-1.bin", static_cast(program)), std::ios::binary); + dump.write(reinterpret_cast(bytes.data()), bytes.size()); + } + } + InitializeNativeCode(); + } gp_passthrough_mask = maxwell3d->regs.post_vtg_shader_attrib_skip_mask; switch (program) { case Maxwell::ShaderType::VertexA: @@ -331,8 +497,7 @@ GraphicsEnvironment::GraphicsEnvironment(Tegra::Engines::Maxwell3D& maxwell3d_, const u64 local_size{sph.LocalMemorySize()}; ASSERT(local_size <= (std::numeric_limits::max)()); local_memory_size = static_cast(local_size) + sph.common3.shader_local_memory_crs_size; - texture_bound = maxwell3d->regs.bindless_texture_const_buffer_slot; - is_proprietary_driver = texture_bound == 2; + is_proprietary_driver = maxwell3d->GetNativeMacroEngine().native_turing || texture_bound == 2; has_hle_engine_state = maxwell3d->engine_state == Tegra::Engines::Maxwell3D::EngineHint::OnHLEMacro; } @@ -417,6 +582,12 @@ ComputeEnvironment::ComputeEnvironment(Tegra::Engines::KeplerCompute& kepler_com is_proprietary_driver = texture_bound == 2; shared_memory_size = qmd.shared_alloc; workgroup_size = {qmd.block_dim_x, qmd.block_dim_y, qmd.block_dim_z}; + if (kepler_compute->IsNativeMode()) { + initial_offset = sizeof(sph); + is_proprietary_driver = true; + InitializeNativeCode(0); + start_address = sizeof(sph); + } } u32 ComputeEnvironment::ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) { @@ -511,7 +682,7 @@ void FileEnvironment::Deserialize(std::ifstream& file) { if (stage == Shader::Stage::Compute) { file.read(reinterpret_cast(&workgroup_size), sizeof(workgroup_size)) .read(reinterpret_cast(&shared_memory_size), sizeof(shared_memory_size)); - initial_offset = 0; + initial_offset = Shader::Ampere::HasNativeAddressMap(code) ? sizeof(sph) : 0; } else { file.read(reinterpret_cast(&sph), sizeof(sph)); initial_offset = sizeof(sph); @@ -526,6 +697,13 @@ void FileEnvironment::Dump(u64 pipeline_hash, u64 shader_hash) { DumpImpl(pipeline_hash, shader_hash, code, read_highest, read_lowest, initial_offset, stage); } +u32 FileEnvironment::TranslateProgramAddress(u32 address) const { + if (const auto translated = Shader::Ampere::NativeProgramAddress(code, address)) return *translated; + if (Shader::Ampere::HasNativeAddressMap(code)) + throw Shader::RuntimeError("Invalid cached native indirect branch address {:#x}", address); + return address; +} + u64 FileEnvironment::ReadInstruction(u32 address) { if (address < read_lowest || address > read_highest) { throw Shader::LogicError("Out of bounds address {}", address); diff --git a/src/video_core/shader_environment.h b/src/video_core/shader_environment.h index 924e61ef46..e49f01d95f 100644 --- a/src/video_core/shader_environment.h +++ b/src/video_core/shader_environment.h @@ -47,6 +47,8 @@ public: [[nodiscard]] u64 ReadInstruction(u32 address) final; + [[nodiscard]] u32 TranslateProgramAddress(u32 address) const override; + [[nodiscard]] std::optional Analyze(); void SetCachedSize(size_t size_bytes); @@ -57,6 +59,10 @@ public: [[nodiscard]] size_t ReadSizeBytes() const noexcept; + [[nodiscard]] size_t SourceSizeBytes() const noexcept { + return native_source_size; + } + [[nodiscard]] bool CanBeSerialized() const noexcept; [[nodiscard]] u64 CalculateHash() const; @@ -72,6 +78,9 @@ public: protected: std::optional TryFindSize(); + void InitializeNativeCode(u32 source_header_size = 128); + size_t native_source_size{}; + Tegra::Texture::TICEntry ReadTextureInfo(GPUVAddr tic_addr, u32 tic_limit, bool via_header_index, u32 raw); @@ -172,6 +181,8 @@ public: [[nodiscard]] u64 ReadInstruction(u32 address) override; + [[nodiscard]] u32 TranslateProgramAddress(u32 address) const override; + [[nodiscard]] u32 ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) override; [[nodiscard]] Shader::TextureType ReadTextureType(u32 handle) override; diff --git a/src/video_core/texture_cache/accelerated_swizzle.cpp b/src/video_core/texture_cache/accelerated_swizzle.cpp index 4c3f724d79..8ae0ed218d 100644 --- a/src/video_core/texture_cache/accelerated_swizzle.cpp +++ b/src/video_core/texture_cache/accelerated_swizzle.cpp @@ -37,6 +37,7 @@ BlockLinearSwizzle2DParams MakeBlockLinearSwizzle2DParams(const SwizzleParameter .x_shift = GOB_SIZE_SHIFT + block.height + block.depth, .block_height = block.height, .block_height_mask = (1U << block.height) - 1, + .native_tiling = info.native_tiling, }; } @@ -63,6 +64,7 @@ BlockLinearSwizzle3DParams MakeBlockLinearSwizzle3DParams(const SwizzleParameter .block_height_mask = (1U << block.height) - 1, .block_depth = block.depth, .block_depth_mask = (1U << block.depth) - 1, + .native_tiling = info.native_tiling, }; } diff --git a/src/video_core/texture_cache/accelerated_swizzle.h b/src/video_core/texture_cache/accelerated_swizzle.h index d4250da689..8296950c38 100644 --- a/src/video_core/texture_cache/accelerated_swizzle.h +++ b/src/video_core/texture_cache/accelerated_swizzle.h @@ -20,6 +20,7 @@ struct BlockLinearSwizzle2DParams { u32 x_shift; u32 block_height; u32 block_height_mask; + u32 native_tiling; }; struct BlockLinearSwizzle3DParams { @@ -33,6 +34,7 @@ struct BlockLinearSwizzle3DParams { u32 block_height_mask; u32 block_depth; u32 block_depth_mask; + u32 native_tiling; }; [[nodiscard]] BlockLinearSwizzle2DParams MakeBlockLinearSwizzle2DParams( diff --git a/src/video_core/texture_cache/format_lookup_table.cpp b/src/video_core/texture_cache/format_lookup_table.cpp index d6e019cddf..20bb96aae3 100644 --- a/src/video_core/texture_cache/format_lookup_table.cpp +++ b/src/video_core/texture_cache/format_lookup_table.cpp @@ -157,6 +157,7 @@ PixelFormat PixelFormatFromTextureInfo(TextureFormat format, ComponentType red, return PixelFormat::S8_UINT_D24_UNORM; case Hash(TextureFormat::G24R8, UINT, UNORM, UNORM, UNORM, LINEAR): return PixelFormat::S8_UINT_D24_UNORM; + case Hash(TextureFormat::S8Z24, UNORM, UINT, UNORM, UNORM, LINEAR): case Hash(TextureFormat::S8Z24, UNORM, UINT, UINT, UINT, LINEAR): return PixelFormat::D24_UNORM_S8_UINT; case Hash(TextureFormat::Z32_X24S8, FLOAT, UINT, UNORM, UNORM, LINEAR): diff --git a/src/video_core/texture_cache/image_info.h b/src/video_core/texture_cache/image_info.h index eb490a6424..7407062d19 100644 --- a/src/video_core/texture_cache/image_info.h +++ b/src/video_core/texture_cache/image_info.h @@ -14,6 +14,17 @@ namespace VideoCommon { using Tegra::Texture::TICEntry; using VideoCore::Surface::PixelFormat; +// Native descriptors address the selected view directly. The legacy NVN +// layer-base bookkeeping is not an additional array prefix for this ABI. +inline TICEntry TextureViewHeader(TICEntry entry, bool native) { + if (native) { + entry.layer_base_0_2.Assign(0); + entry.layer_base_3_7.Assign(0); + entry.layer_base_8_10.Assign(0); + } + return entry; +} + struct ImageInfo { ImageInfo() = default; explicit ImageInfo(const TICEntry& config) noexcept; @@ -37,6 +48,7 @@ struct ImageInfo { u32 maybe_unaligned_layer_stride = 0; u32 num_samples = 1; u32 tile_width_spacing = 0; + bool native_tiling = false; bool rescaleable = false; bool downscaleable = false; bool forced_flushed = false; diff --git a/src/video_core/texture_cache/pending_download.h b/src/video_core/texture_cache/pending_download.h new file mode 100644 index 0000000000..52a48e82d3 --- /dev/null +++ b/src/video_core/texture_cache/pending_download.h @@ -0,0 +1,58 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later +#pragma once + +#include +#include +#include +#include +#include +#include "common/slot_vector.h" + +namespace VideoCommon { + +template +struct PendingImageDownload { + bool is_swizzle; + size_t async_buffer_id; + Common::SlotId object_id; + std::shared_ptr retired_image{}; + + Image& GetImage(Common::SlotVector& images) const { + return retired_image ? *retired_image : images[object_id]; + } +}; + +// Whole-image copies are recorded together at commit time. Repeated requests +// for the same live image therefore observe the same contents. Keep the first +// occurrence so reverse CPU writeback order remains unchanged for aliases. +template +size_t DeduplicatePendingImageDownloads(Downloads& downloads, Common::SlotVector& images) { + std::unordered_set seen; + const size_t previous_size = downloads.size(); + const auto end = std::remove_if(downloads.begin(), downloads.end(), [&](const auto& download) { + return download.is_swizzle && !seen.insert(&download.GetImage(images)).second; + }); + downloads.erase(end, downloads.end()); + return previous_size - downloads.size(); +} + +// Cache slots may be recycled before queued GPU downloads are submitted or +// consumed. Keep the actual image, including its metadata, until those finish. +template +bool RetainPendingImageDownloads(Common::SlotId id, Image& image, Downloads& pending, + Batches& submitted) { + std::shared_ptr retained; + const auto retain = [&](auto& downloads) { + for (auto& download : downloads) { + if (!download.is_swizzle || download.object_id != id || download.retired_image) continue; + if (!retained) retained = std::make_shared(std::move(image)); + download.retired_image = retained; + } + }; + retain(pending); + for (auto& downloads : submitted) retain(downloads); + return retained != nullptr; +} + +} // namespace VideoCommon diff --git a/src/video_core/texture_cache/texture_cache.h b/src/video_core/texture_cache/texture_cache.h index e92e54f8ce..cd4ef82dbe 100644 --- a/src/video_core/texture_cache/texture_cache.h +++ b/src/video_core/texture_cache/texture_cache.h @@ -803,6 +803,17 @@ std::pair TextureCache

::TryFindFramebufferImage if (it == page_table.end()) { return {}; } + const auto view_format = [&]() { + switch (config.pixel_format) { + case Service::android::PixelFormat::Rgb565: + return PixelFormat::R5G6B5_UNORM; + case Service::android::PixelFormat::Bgra8888: + return PixelFormat::B8G8R8A8_UNORM; + default: + return PixelFormat::A8B8G8R8_UNORM; + } + }(); + const auto& image_map_ids = it->second; boost::container::small_vector valid_image_ids; for (const ImageMapId map_id : image_map_ids) { @@ -814,20 +825,17 @@ std::pair TextureCache

::TryFindFramebufferImage if (image.image_view_ids.empty()) { continue; } + if (!VideoCore::Surface::IsViewCompatible(image.info.format, view_format, false, true)) { + static u32 rejected_alias_count{}; + if (std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS") && rejected_alias_count++ < 8) { + LOG_WARNING(HW_GPU, "Ignoring incompatible framebuffer alias: address={:#x}, image={}, view={}", + cpu_addr, image.info.format, view_format); + } + continue; + } valid_image_ids.push_back(map.image_id); } - const auto view_format = [&]() { - switch (config.pixel_format) { - case Service::android::PixelFormat::Rgb565: - return PixelFormat::R5G6B5_UNORM; - case Service::android::PixelFormat::Bgra8888: - return PixelFormat::B8G8R8A8_UNORM; - default: - return PixelFormat::A8B8G8R8_UNORM; - } - }(); - const auto GetImageViewForFramebuffer = [&](ImageId image_id) { ImageViewInfo info{ImageViewType::e2D, view_format}; if (config.blending == Tegra::BlendMode::Opaque) { @@ -869,6 +877,7 @@ void TextureCache

::CommitAsyncFlushes() { // This is intentionally passing the value by copy if constexpr (IMPLEMENTS_ASYNC_DOWNLOADS) { auto& download_ids = uncommitted_downloads; + const size_t redundant_images = DeduplicatePendingImageDownloads(download_ids, slot_images); if (download_ids.empty()) { committed_downloads.emplace_back(std::move(uncommitted_downloads)); uncommitted_downloads.clear(); @@ -882,17 +891,32 @@ void TextureCache

::CommitAsyncFlushes() { for (PendingDownload& download_info : download_ids) { if (download_info.is_swizzle) { total_size_bytes += - Common::AlignUp(slot_images[download_info.object_id].unswizzled_size_bytes, 64); + Common::AlignUp(download_info.GetImage(slot_images).unswizzled_size_bytes, 64); any_none_dma = true; download_info.async_buffer_id = last_async_buffer_id; } } if (any_none_dma) { + if (std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS") && total_size_bytes >= 64_MiB) { + LOG_WARNING(HW_GPU, "Large image readback batch: bytes={}, entries={}, duplicates_removed={}, submitted_batches={}, retained_buffers={}", + total_size_bytes, download_ids.size(), redundant_images, + committed_downloads.size(), async_buffers_death_ring.size()); + size_t traced_images = 0; + for (const auto& download : download_ids) { + if (!download.is_swizzle || traced_images++ >= 8) continue; + const auto& image = download.GetImage(slot_images); + LOG_WARNING(HW_GPU, "Readback image: gpu={:#x}, cpu={:#x}, format={}, size={}x{}x{}, levels={}, layers={}, layer_stride={}, guest_bytes={}, download_bytes={}", + image.gpu_addr, image.cpu_addr, image.info.format, + image.info.size.width, image.info.size.height, image.info.size.depth, + image.info.resources.levels, image.info.resources.layers, + image.info.layer_stride, image.guest_size_bytes, image.unswizzled_size_bytes); + } + } auto download_map = runtime.DownloadStagingBuffer(total_size_bytes, true); for (const PendingDownload& download_info : download_ids) { if (download_info.is_swizzle) { - Image& image = slot_images[download_info.object_id]; + Image& image = download_info.GetImage(slot_images); const auto copies = FullDownloadCopies(image.info); image.DownloadMemory(download_map, copies); download_map.offset += Common::AlignUp(image.unswizzled_size_bytes, 64); @@ -925,7 +949,7 @@ void TextureCache

::PopAsyncFlushes() { auto& download_info = download_ids[i - 1]; auto& download_buffer = download_map[download_info.async_buffer_id]; if (download_info.is_swizzle) { - const ImageBase& image = slot_images[download_info.object_id]; + const ImageBase& image = download_info.GetImage(slot_images); const auto copies = FullDownloadCopies(image.info); download_buffer.offset -= Common::AlignUp(image.unswizzled_size_bytes, 64); std::span download_span = @@ -954,7 +978,7 @@ void TextureCache

::PopAsyncFlushes() { size_t total_size_bytes = 0; for (const PendingDownload& download_info : download_ids) { if (download_info.is_swizzle) { - total_size_bytes += slot_images[download_info.object_id].unswizzled_size_bytes; + total_size_bytes += download_info.GetImage(slot_images).unswizzled_size_bytes; } } auto download_map = runtime.DownloadStagingBuffer(total_size_bytes); @@ -963,7 +987,7 @@ void TextureCache

::PopAsyncFlushes() { if (!download_info.is_swizzle) { continue; } - Image& image = slot_images[download_info.object_id]; + Image& image = download_info.GetImage(slot_images); const auto copies = FullDownloadCopies(image.info); image.DownloadMemory(download_map, copies); download_map.offset += image.unswizzled_size_bytes; @@ -976,7 +1000,7 @@ void TextureCache

::PopAsyncFlushes() { if (!download_info.is_swizzle) { continue; } - const ImageBase& image = slot_images[download_info.object_id]; + const ImageBase& image = download_info.GetImage(slot_images); const auto copies = FullDownloadCopies(image.info); SwizzleImage(*gpu_memory, image.gpu_addr, image.info, copies, download_span, swizzle_data_buffer); download_map.offset += image.unswizzled_size_bytes; @@ -1178,7 +1202,18 @@ void TextureCache

::UploadImageContents(Image& image, StagingBuffer& staging) } template -ImageViewId TextureCache

::CreateImageView(const TICEntry& config) { +ImageViewId TextureCache

::CreateImageView(const TICEntry& raw_config) { + const bool native = gpu_memory->UsesNativeTiling(); + if (native && raw_config.BaseLayer() && std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS")) { + static u32 texture_view_traces{}; + if (texture_view_traces++ < 8) { + LOG_WARNING(HW_GPU, "Native texture view header: address={:#x}, dimensions={}x{}x{}, type={}, legacy_layer={}, raw={:016x}/{:016x}/{:016x}/{:016x}", + raw_config.Address(), raw_config.Width(), raw_config.Height(), raw_config.Depth(), + static_cast(raw_config.texture_type.Value()), raw_config.BaseLayer(), + raw_config.raw[0], raw_config.raw[1], raw_config.raw[2], raw_config.raw[3]); + } + } + const TICEntry config = TextureViewHeader(raw_config, native); const ImageInfo info(config); if (info.type == ImageType::Buffer) { const ImageViewInfo view_info(config, 0); @@ -1526,7 +1561,9 @@ ImageId TextureCache

::InsertImage(const ImageInfo& info, GPUVAddr gpu_addr, } } ASSERT_MSG(cpu_addr, "Tried to insert an image to an invalid gpu_addr={:#x}", gpu_addr); - const ImageId image_id = JoinImages(info, gpu_addr, *cpu_addr); + ImageInfo storage_info = info; + storage_info.native_tiling = gpu_memory->UsesNativeTiling(); + const ImageId image_id = JoinImages(storage_info, gpu_addr, *cpu_addr); const Image& image = slot_images[image_id]; // Using "image.gpu_addr" instead of "gpu_addr" is important because it might be different const auto [it, is_new] = image_allocs_table.try_emplace(image.gpu_addr); @@ -2385,7 +2422,11 @@ void TextureCache

::DeleteImage(ImageId image_id, bool immediate_delete) { } slot_image_views.erase(image_view_id); } - if (!immediate_delete) { + const bool retained = RetainPendingImageDownloads(image_id, slot_images[image_id], + uncommitted_downloads, committed_downloads); + if (retained) { + LOG_INFO(Render_Vulkan, "Retained image for queued download: id={}", image_id.Value()); + } else if (!immediate_delete) { sentenced_images.Push(std::move(slot_images[image_id])); } slot_images.erase(image_id); diff --git a/src/video_core/texture_cache/texture_cache_base.h b/src/video_core/texture_cache/texture_cache_base.h index 8a2318cc72..d103d66189 100644 --- a/src/video_core/texture_cache/texture_cache_base.h +++ b/src/video_core/texture_cache/texture_cache_base.h @@ -28,6 +28,7 @@ #include #include "common/scratch_buffer.h" #include "common/slot_vector.h" +#include "video_core/texture_cache/pending_download.h" #include "common/thread_worker.h" #include "video_core/compatible_formats.h" #include "video_core/control/channel_state_cache.h" @@ -463,11 +464,7 @@ private: size_t size; }; - struct PendingDownload { - bool is_swizzle; - size_t async_buffer_id; - Common::SlotId object_id; - }; + using PendingDownload = PendingImageDownload; Common::SlotVector slot_images; Common::SlotVector slot_map_views; diff --git a/src/video_core/texture_cache/util.cpp b/src/video_core/texture_cache/util.cpp index e55d0752ec..21070b2d25 100644 --- a/src/video_core/texture_cache/util.cpp +++ b/src/video_core/texture_cache/util.cpp @@ -439,7 +439,7 @@ template } layers = 1; } else { - layers = (std::max)(resources.layers, info.resources.layers + base->layer); + layers = (std::max)(info.resources.layers, resources.layers + base->layer); } return OverlapResult{ .gpu_addr = overlap.gpu_addr, @@ -566,7 +566,7 @@ void SwizzleBlockLinearImage(Tegra::MemoryManager& gpu_memory, GPUVAddr gpu_addr dst(gpu_memory, gpu_addr + guest_offset, subresource_size, &tmp_buffer); SwizzleTexture(dst, src, bytes_per_block, num_tiles.width, num_tiles.height, - num_tiles.depth, block.height, block.depth); + num_tiles.depth, block.height, block.depth, 1, gpu_memory.UsesNativeTiling()); } host_offset += host_bytes_per_layer; @@ -913,7 +913,7 @@ boost::container::small_vector UnswizzleImage(Tegra::Memory const std::span dst = output.subspan(host_offset); const std::span src = input.subspan(guest_offset + guest_layer_offset); UnswizzleTexture(dst, src, 1U << bpp_log2, num_tiles.width, num_tiles.height, - num_tiles.depth, block.height, block.depth, stride_alignment); + num_tiles.depth, block.height, block.depth, stride_alignment, gpu_memory.UsesNativeTiling()); guest_layer_offset += layer_stride; host_offset += host_bytes_per_layer; } -- cgit v1.2.3 From c01459347405591446d50858d7b075d617a58be9 Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Thu, 29 Oct 2026 01:26:52 -0400 Subject: video_core: cheaper small GPU reads, buffer usage tracking and compute uniforms - MemoryManager::ReadBlockUnsafe reads within one 4 KiB page through a single translation (as Read does) instead of the page walk: texture and sampler descriptors, several per draw (2.6% -> 0.1% of title 1's GPU thread). - UsageTracker skips ranges it tracked just before (until Reset): draws bind the same multi-megabyte storage buffers again and again, and each Track wrote a word per 4 KiB (2.0% -> 0.7%). - Small compute uniform buffers the GPU has not written are streamed, as graphics ones are; synchronizing them re-protected pages an inline upload had just unprotected, two VirtualProtect calls per dispatch (1.9% -> 0.7%, page protection 3.2% -> 2.3%). --- src/video_core/buffer_cache/buffer_cache.h | 9 ++++++++- src/video_core/buffer_cache/usage_tracker.h | 17 +++++++++++++++++ src/video_core/memory_manager.cpp | 10 ++++++++++ 3 files changed, 35 insertions(+), 1 deletion(-) (limited to 'src/video_core/memory_manager.cpp') diff --git a/src/video_core/buffer_cache/buffer_cache.h b/src/video_core/buffer_cache/buffer_cache.h index 4df855c0fe..ad2e200438 100644 --- a/src/video_core/buffer_cache/buffer_cache.h +++ b/src/video_core/buffer_cache/buffer_cache.h @@ -1211,8 +1211,15 @@ void BufferCache

::BindHostComputeUniformBuffers() { return alignment > 1 && (offset % alignment) != 0; } }(); + // Small buffers the GPU has not written are streamed, as graphics ones are: synchronizing + // them made every dispatch re-protect pages an inline upload had just unprotected (two + // VirtualProtect calls per dispatch, ~4,000 dispatches a second in title 1). + const bool use_fast_buffer = + needs_alignment_stream || + (has_host_buffer && size <= channel_state->uniform_buffer_skip_cache_size && + !memory_tracker.IsRegionGpuModified(binding.device_addr, size)); if constexpr (!IS_OPENGL) { - if (needs_alignment_stream) { + if (use_fast_buffer) { const std::span span = runtime.BindMappedUniformBuffer(0, binding_index, size); device_memory.ReadBlockUnsafe(binding.device_addr, span.data(), size); diff --git a/src/video_core/buffer_cache/usage_tracker.h b/src/video_core/buffer_cache/usage_tracker.h index 6f84e463e3..3e2da8bc0b 100644 --- a/src/video_core/buffer_cache/usage_tracker.h +++ b/src/video_core/buffer_cache/usage_tracker.h @@ -6,6 +6,10 @@ #pragma once +#include +#include +#include + #include "common/alignment.h" #include "common/common_types.h" @@ -24,9 +28,19 @@ public: void Reset() noexcept { std::ranges::fill(pages, 0ULL); + recent.fill({~u64{0}, 0}); } void Track(u64 offset, u64 size) noexcept { + // Tracking is idempotent until Reset: skip ranges tracked just before. Draws bind the + // same (often multi-megabyte) storage buffers again and again, and each Track wrote a + // word per 4 KiB of the range. + for (const auto& [recent_offset, recent_size] : recent) { + if (recent_offset == offset && recent_size == size) { + return; + } + } + recent[recent_cursor++ % recent.size()] = {offset, size}; const size_t page = offset >> BUFFER_PAGE_SHIFT; const size_t page_end = (offset + size) >> BUFFER_PAGE_SHIFT; if (page_end < page || page_end >= pages.size()) { @@ -83,6 +97,9 @@ private: private: std::vector pages; + /// Ranges tracked last, (offset, size). + std::array, 4> recent{{{~u64{0}, 0}, {~u64{0}, 0}, {~u64{0}, 0}, {~u64{0}, 0}}}; + size_t recent_cursor{}; }; } // namespace VideoCommon diff --git a/src/video_core/memory_manager.cpp b/src/video_core/memory_manager.cpp index 9a9fdab7bc..0eae8fac24 100644 --- a/src/video_core/memory_manager.cpp +++ b/src/video_core/memory_manager.cpp @@ -410,6 +410,16 @@ void MemoryManager::ReadBlock(GPUVAddr gpu_src_addr, void* dest_buffer, std::siz } void MemoryManager::ReadBlockUnsafe(GPUVAddr gpu_src_addr, void* dest_buffer, const std::size_t size) const { + // Small reads within one 4 KiB page (texture and sampler descriptors, several per draw): one + // translation and a copy, as Read does, instead of the page walk (~3% of the GPU thread + // in title 1). Anything else, unmapped or sparse pages included, takes the full path. + constexpr std::size_t SmallPage = 0x1000; + if (size <= SmallPage && (gpu_src_addr & (SmallPage - 1)) + size <= SmallPage) { + if (const u8* const pointer = GetPointer(gpu_src_addr)) { + std::memcpy(dest_buffer, pointer, size); + return; + } + } ReadBlockImpl(gpu_src_addr, dest_buffer, size, VideoCommon::CacheType::None, true); } -- cgit v1.2.3